/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgesdd (SVD with Divide & Conquer)
Please keep code in testing_zgesdd.cpp and testing_zgesvd.cpp similar.
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
magmaDoubleComplex *h_A, *h_R, *U, *VT, *h_work;
magmaDoubleComplex dummy[1];
double *S1, *S2;
#ifdef COMPLEX
magma_int_t lrwork=0;
double *rwork;
#endif
magma_int_t *iwork;
magma_int_t M, N, N_U, M_VT, lda, ldu, ldv, n2, min_mn, max_mn, info, nb, lwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_vec_t jobz;
magma_int_t status = 0;
MAGMA_UNUSED( max_mn ); // used only in complex
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
jobz = opts.jobu;
magma_vec_t jobs[] = { MagmaNoVec, MagmaSomeVec, MagmaOverwriteVec, MagmaAllVec };
if ( opts.check && ! opts.all && (jobz == MagmaNoVec)) {
printf( "%% NOTE: some checks require that singular vectors are computed;\n"
"%% set jobz (option -U[NASO]) to be S, O, or A.\n\n" );
}
printf("%% jobz M N CPU time (sec) GPU time (sec) |S1-S2| |A-USV^H| |I-UU^H|/M |I-VV^H|/N S sorted\n");
printf("%%==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int ijobz = 0; ijobz < 4; ++ijobz ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
if ( opts.all ) {
jobz = jobs[ ijobz ];
}
else if ( ijobz > 0 ) {
// if not testing all, run only once, when ijobz = 0,
// but jobz come from opts (above loops).
continue;
}
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
max_mn = max(M, N);
N_U = (jobz == MagmaAllVec ? M : min_mn);
M_VT = (jobz == MagmaAllVec ? N : min_mn);
lda = M;
ldu = M;
ldv = M_VT;
n2 = lda*N;
nb = magma_get_zgesvd_nb( M, N );
// x and y abbreviations used in zgesdd and dgesdd documentation
magma_int_t x = max(M,N);
magma_int_t y = min(M,N);
#ifdef COMPLEX
bool tall = (x >= int(y*17/9.)); // true if tall (m >> n) or wide (n >> m)
#else
bool tall = (x >= int(y*11/6.)); // true if tall (m >> n) or wide (n >> m)
#endif
// query or use formula for workspace size
switch( opts.svd_work ) {
case 0: { // query for workspace size
lwork = -1;
magma_zgesdd( jobz, M, N,
NULL, lda, NULL, NULL, ldu, NULL, ldv, dummy, lwork,
#ifdef COMPLEX
NULL,
#endif
NULL, &info );
lwork = (int) MAGMA_Z_REAL( dummy[0] );
break;
}
case 1: // minimum
case 2: // optimal
case 3: { // optimal (for gesdd, 2 & 3 are same; for gesvd, they differ)
// formulas from zgesdd and dgesdd documentation
bool sml = (opts.svd_work == 1); // 1 is small workspace, 2,3 are large workspace
#ifdef COMPLEX // ----------------------------------------
if (jobz == MagmaNoVec) {
if (tall) { lwork = 2*y + (2*y)*nb; }
else { lwork = 2*y + (x+y)*nb; }
}
if (jobz == MagmaOverwriteVec) {
if (tall) {
if (sml) { lwork = 2*y*y + 2*y + (2*y)*nb; }
else { lwork = y*y + x*y + 2*y + (2*y)*nb; } // not big deal
}
else {
//if (sml) { lwork = 2*y + max( (x+y)*nb, y*y + y ); }
//else { lwork = 2*y + max( (x+y)*nb, x*y + y*nb ); }
// LAPACK 3.4.2 over-estimates workspaces. For compatability, use these:
if (sml) { lwork = 2*y + max( (x+y)*nb, y*y + x ); }
else { lwork = 2*y + (x+y)*nb + x*y; }
}
}
if (jobz == MagmaSomeVec) {
if (tall) { lwork = y*y + 2*y + (2*y)*nb; }
else { lwork = 2*y + (x+y)*nb; }
}
if (jobz == MagmaAllVec) {
if (tall) {
if (sml) { lwork = y*y + 2*y + max( (2*y)*nb, x ); }
else { lwork = y*y + 2*y + max( (2*y)*nb, x*nb ); }
}
else { lwork = 2*y + (x+y)*nb; }
}
#else // REAL ----------------------------------------
if (jobz == MagmaNoVec) {
if (tall) { lwork = 3*y + max( (2*y)*nb, 7*y ); }
else { lwork = 3*y + max( (x+y)*nb, 7*y ); }
}
if (jobz == MagmaOverwriteVec) {
if (tall) {
if (sml) { lwork = y*y + 3*y + max( (2*y)*nb, 4*y*y + 4*y ); }
else { lwork = y*y + 3*y + max( max( (2*y)*nb, 4*y*y + 4*y ), y*y + y*nb ); }
}
else {
if (sml) { lwork = 3*y + max( (x+y)*nb, 4*y*y + 4*y ); }
else { lwork = 3*y + max( (x+y)*nb, 3*y*y + 4*y + x*y ); } // extra space not too important?
}
}
if (jobz == MagmaSomeVec) {
if (tall) { lwork = y*y + 3*y + max( (2*y)*nb, 3*y*y + 4*y ); }
else { lwork = 3*y + max( (x+y)*nb, 3*y*y + 4*y ); }
}
if (jobz == MagmaAllVec) {
if (tall) {
if (sml) { lwork = y*y + max( 3*y + max( (2*y)*nb, 3*y*y + 4*y ), y + x ); }
else { lwork = y*y + max( 3*y + max( (2*y)*nb, 3*y*y + 4*y ), y + x*nb ); }
// LAPACK 3.4.2 over-estimates workspaces. For compatability, use these:
//if (sml) { lwork = y*y + 3*y + max( (2*y)*nb, 3*y*y + 3*y + x ); }
//else { lwork = y*y + max( 3*y + max( (2*y)*nb, max( 3*y*y + 3*y + x, 3*y*y + 4*y )), y + x*nb ); }
}
else { lwork = 3*y + max( (x+y)*nb, 3*y*y + 4*y ); }
}
#endif
break;
}
default: {
fprintf( stderr, "Error: unknown option svd_work %d\n", (int) opts.svd_work );
return -1;
break;
}
}
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( VT, magmaDoubleComplex, ldv*N ); // N x N (jobz=A) or min(M,N) x N
TESTING_MALLOC_CPU( U, magmaDoubleComplex, ldu*N_U ); // M x M (jobz=A) or M x min(M,N)
TESTING_MALLOC_CPU( S1, double, min_mn );
TESTING_MALLOC_CPU( S2, double, min_mn );
TESTING_MALLOC_CPU( iwork, magma_int_t, 8*min_mn );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
#ifdef COMPLEX
if (jobz == MagmaNoVec) {
// requires 5*min_mn, but MKL (11.1) seems to have a bug
// requiring 7*min_mn in some cases (e.g., jobz=N, m=100, n=170)
lrwork = 7*min_mn;
}
else if (tall) {
lrwork = 5*min_mn*min_mn + 5*min_mn;
}
else {
lrwork = max( 5*min_mn*min_mn + 5*min_mn,
2*max_mn*min_mn + 2*min_mn*min_mn + min_mn );
}
TESTING_MALLOC_CPU( rwork, double, lrwork );
#endif
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgesdd( jobz, M, N,
h_R, lda, S1, U, ldu, VT, ldv, h_work, lwork,
#ifdef COMPLEX
rwork,
#endif
iwork, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_zgesdd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
double eps = lapackf77_dlamch( "E" );
double result[5] = { -1/eps, -1/eps, -1/eps, -1/eps, -1/eps };
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zcds]drvbd routine.
A is factored as A = U diag(S) VT and the following 4 tests computed:
(1) | A - U diag(S) VT | / ( |A| max(M,N) )
(2) | I - U^H U | / ( M )
(3) | I - VT VT^H | / ( N )
(4) S contains MNMIN nonnegative values in decreasing order.
(Return 0 if true, 1/ULP if false.)
=================================================================== */
magma_int_t izero = 0;
// get size and location of U and V^T depending on jobz
// U2=NULL and VT2=NULL if they were not computed (e.g., jobz=N)
magmaDoubleComplex *U2 = NULL;
magmaDoubleComplex *VT2 = NULL;
if ( jobz == MagmaSomeVec || jobz == MagmaAllVec ) {
U2 = U;
VT2 = VT;
}
else if ( jobz == MagmaOverwriteVec ) {
if ( M >= N ) {
U2 = h_R;
ldu = lda;
VT2 = VT;
}
else {
U2 = U;
VT2 = h_R;
ldv = lda;
}
}
// zbdt01 needs M+N
// zunt01 prefers N*(N+1) to check U; M*(M+1) to check V
magma_int_t lwork_err = M+N;
if ( U2 != NULL ) {
lwork_err = max( lwork_err, N_U*(N_U+1) );
}
if ( VT2 != NULL ) {
lwork_err = max( lwork_err, M_VT*(M_VT+1) );
}
magmaDoubleComplex *h_work_err;
TESTING_MALLOC_CPU( h_work_err, magmaDoubleComplex, lwork_err );
// zbdt01 and zunt01 need max(M,N), depending
double *rwork_err;
TESTING_MALLOC_CPU( rwork_err, double, max(M,N) );
if ( U2 != NULL && VT2 != NULL ) {
// since KD=0 (3rd arg), E is not referenced so pass NULL (9th arg)
lapackf77_zbdt01(&M, &N, &izero, h_A, &lda,
U2, &ldu, S1, NULL, VT2, &ldv,
h_work_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[0]);
}
if ( U2 != NULL ) {
lapackf77_zunt01("Columns", &M, &N_U, U2, &ldu, h_work_err, &lwork_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[1]);
}
if ( VT2 != NULL ) {
lapackf77_zunt01("Rows", &M_VT, &N, VT2, &ldv, h_work_err, &lwork_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[2]);
}
result[3] = 0.;
for (int j=0; j < min_mn-1; j++) {
if ( S1[j] < S1[j+1] )
result[3] = 1.;
if ( S1[j] < 0. )
result[3] = 1.;
}
if (min_mn > 1 && S1[min_mn-1] < 0.)
result[3] = 1.;
result[0] *= eps;
result[1] *= eps;
result[2] *= eps;
TESTING_FREE_CPU( h_work_err );
TESTING_FREE_CPU( rwork_err );
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zgesdd( lapack_vec_const(jobz), &M, &N,
h_A, &lda, S2, U, &ldu, VT, &ldv, h_work, &lwork,
#ifdef COMPLEX
rwork,
#endif
iwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info != 0) {
printf("lapackf77_zgesdd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
double work[1], c_neg_one = -1;
blasf77_daxpy(&min_mn, &c_neg_one, S1, &ione, S2, &ione);
result[4] = lapackf77_dlange("f", &min_mn, &ione, S2, &min_mn, work);
result[4] /= lapackf77_dlange("f", &min_mn, &ione, S1, &min_mn, work);
printf(" %c %5d %5d %7.2f %7.2f %8.2e",
lapack_vec_const(jobz)[0],
(int) M, (int) N, cpu_time, gpu_time, result[4] );
}
else {
printf(" %c %5d %5d --- %7.2f --- ",
lapack_vec_const(jobz)[0],
(int) M, (int) N, gpu_time );
}
if ( opts.check ) {
if ( result[0] < 0. ) { printf(" --- "); } else { printf(" %#9.3g", result[0]); }
if ( result[1] < 0. ) { printf(" --- "); } else { printf(" %#9.3g", result[1]); }
if ( result[2] < 0. ) { printf(" --- "); } else { printf(" %#9.3g", result[2]); }
bool okay = (result[0] < tol) && (result[1] < tol) && (result[2] < tol) && (result[3] == 0.) && (result[4] < tol);
printf(" %3s %s\n", (result[3] == 0. ? "yes" : "no"), (okay ? "ok" : "failed"));
status += ! okay;
}
else {
printf("\n");
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( VT );
TESTING_FREE_CPU( U );
TESTING_FREE_CPU( S1 );
TESTING_FREE_CPU( S2 );
TESTING_FREE_CPU( iwork );
#ifdef COMPLEX
TESTING_FREE_CPU( rwork );
#endif
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
fflush( stdout );
}}
if ( opts.all || opts.niter > 1 ) {
printf("\n");
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zungqr_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *hA, *hR, *tau, *h_work;
magmaDoubleComplex *dA, *dT;
magma_int_t m, n, k;
magma_int_t n2, lda, ldda, lwork, min_mn, nb, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" m n k CPU GFlop/s (sec) GPU GFlop/s (sec) ||R|| / ||A||\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];
k = opts.ksize[i];
if ( m < n || n < k ) {
printf( "skipping m %d, n %d, k %d because m < n or n < k\n", (int) m, (int) n, (int) k );
continue;
}
lda = m;
ldda = ((m + 31)/32)*32;
n2 = lda*n;
min_mn = min(m, n);
nb = magma_get_zgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_ZUNGQR( m, n, k ) / 1e9;
TESTING_MALLOC_PIN( hA, magmaDoubleComplex, lda*n );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( hR, magmaDoubleComplex, lda*n );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*n );
TESTING_MALLOC_DEV( dT, magmaDoubleComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_zlarnv( &ione, ISEED, &n2, hA );
lapackf77_zlacpy( MagmaUpperLowerStr, &m, &n, hA, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( m, n, hA, lda, dA, ldda );
magma_zgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_zgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
gpu_time = magma_wtime();
magma_zungqr_gpu( m, n, k, dA, ldda, tau, dT, nb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get dA back to the CPU to compare with the CPU result.
magma_zgetmatrix( m, n, dA, ldda, hR, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
error = lapackf77_zlange("f", &m, &n, hA, &lda, work );
lapackf77_zgeqrf( &m, &n, hA, &lda, tau, h_work, &lwork, &info );
if (info != 0)
printf("lapackf77_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
cpu_time = magma_wtime();
lapackf77_zungqr( &m, &n, &k, hA, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_zaxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_zlange("f", &m, &n, hR, &lda, work) / error;
printf("%5d %5d %5d %7.1f (%7.2f) %7.1f (%7.2f) %8.2e\n",
(int) m, (int) n, (int) k,
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
}
else {
printf("%5d %5d %5d --- ( --- ) %7.1f (%7.2f) --- \n",
(int) m, (int) n, (int) k,
gpu_perf, gpu_time );
}
TESTING_FREE_PIN( hA );
TESTING_FREE_PIN( h_work );
TESTING_FREE_CPU( hR );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
// --------------------
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_time=0, cpu_perf=0, gpu_time, gpu_perf, mgpu_time, mgpu_perf, cuda_time, cuda_perf;
double error=0, error2=0, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t n_local[MagmaMaxGPUs];
magma_int_t N, Noffset, lda, ldda, blocks, lhwork, ldwork, matsize, vecsize;
magma_int_t incx = 1;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ylapack, *Ycublas, *Ymagma, *Ymagma1, *hwork;
magmaDoubleComplex_ptr dA, dX, dY;
magmaDoubleComplex_ptr d_lA[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magma_device_t dev;
magma_queue_t queues[MagmaMaxGPUs];
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
magma_int_t nb = 64; // required by magmablas_zhemv_mgpu implementation
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_create( &queues[dev] );
}
// currently, tests all offsets in the offsets array;
// comment out loop below to test a specific offset.
magma_int_t offset = opts.offset;
magma_int_t offsets[] = { 0, 1, 31, 32, 33, 63, 64, 65, 100, 200 };
magma_int_t noffsets = sizeof(offsets) / sizeof(*offsets);
printf("uplo = %s, ngpu %d, block size = %d, offset %d\n",
lapack_uplo_const(opts.uplo), (int) opts.ngpu, (int) nb, (int) offset );
printf( " BLAS CUBLAS MAGMA 1 GPU MAGMA MGPU Error rel Error rel\n"
" N offset Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) to CUBLAS to LAPACK\n"
"===================================================================================================================\n" );
for( int itest = 0; itest < opts.ntest; ++itest ) {
// comment out these two lines & end of loop to test a specific offset
for( int ioffset=0; ioffset < noffsets; ioffset += 1 ) {
offset = offsets[ioffset];
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
Noffset = N + offset;
lda = Noffset;
ldda = ((Noffset+31)/32)*32;
matsize = Noffset*ldda;
vecsize = (Noffset-1)*incx + 1;
gflops = FLOPS_ZHEMV( N ) / 1e9;
blocks = (N + (offset % nb) - 1)/nb + 1;
lhwork = N*opts.ngpu;
ldwork = ldda*(blocks + 1);
TESTING_MALLOC_CPU( A, magmaDoubleComplex, matsize );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ycublas, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma1, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ylapack, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( X, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( hwork, magmaDoubleComplex, lhwork );
magma_setdevice( opts.device );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, matsize );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, vecsize );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, vecsize );
// TODO make magma_zmalloc_bcyclic helper function?
for( dev=0; dev < opts.ngpu; dev++ ) {
n_local[dev] = ((Noffset/nb)/opts.ngpu)*nb;
if (dev < (Noffset/nb) % opts.ngpu)
n_local[dev] += nb;
else if (dev == (Noffset/nb) % opts.ngpu)
n_local[dev] += Noffset % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaDoubleComplex, ldda*n_local[dev] );
TESTING_MALLOC_DEV( dwork[dev], magmaDoubleComplex, ldwork );
}
//////////////////////////////////////////////////////////////////////////
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &matsize, A );
magma_zmake_hermitian( Noffset, A, lda );
lapackf77_zlarnv( &ione, ISEED, &vecsize, X );
lapackf77_zlarnv( &ione, ISEED, &vecsize, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_setdevice( opts.device );
magma_zsetmatrix( Noffset, Noffset, A, lda, dA, ldda );
magma_zsetvector( Noffset, X, incx, dX, incx );
magma_zsetvector( Noffset, Y, incx, dY, incx );
cuda_time = magma_sync_wtime(0);
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo), N,
&alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
&beta, dY + offset, incx );
cuda_time = magma_sync_wtime(0) - cuda_time;
cuda_perf = gflops / cuda_time;
magma_zgetvector( Noffset, dY, incx, Ycublas, incx );
/* =====================================================================
Performs operation using MAGMABLAS (1 GPU)
=================================================================== */
magma_setdevice( opts.device );
magma_zsetvector( Noffset, Y, incx, dY, incx );
gpu_time = magma_sync_wtime( opts.queue );
magmablas_zhemv_work( opts.uplo, N,
alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
beta, dY + offset, incx, dwork[ opts.device ], ldwork,
opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_zgetvector( Noffset, dY, incx, Ymagma1, incx );
/* =====================================================================
Performs operation using MAGMABLAS (multi-GPU)
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( Noffset, Noffset, A, lda, d_lA, ldda, opts.ngpu, nb );
blasf77_zcopy( &Noffset, Y, &incx, Ymagma, &incx );
// workspaces do NOT need to be zero -- set to NAN to prove
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magmablas_zlaset( MagmaFull, ldwork, 1, MAGMA_Z_NAN, MAGMA_Z_NAN, dwork[dev], ldwork );
}
lapackf77_zlaset( "Full", &lhwork, &ione, &MAGMA_Z_NAN, &MAGMA_Z_NAN, hwork, &lhwork );
mgpu_time = magma_sync_wtime(0);
magma_int_t info;
info = magmablas_zhemv_mgpu(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if ( info != 0 )
printf("magmablas_zhemv_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
info = magmablas_zhemv_mgpu_sync(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if ( info != 0 )
printf("magmablas_zhemv_sync returned error %d: %s.\n",
(int) info, magma_strerror( info ));
mgpu_time = magma_sync_wtime(0) - mgpu_time;
mgpu_perf = gflops / mgpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
blasf77_zcopy( &Noffset, Y, &incx, Ylapack, &incx );
cpu_time = magma_wtime();
blasf77_zhemv( lapack_uplo_const(opts.uplo), &N,
&alpha, A + offset + offset*lda, &lda,
X + offset, &incx,
&beta, Ylapack + offset, &incx );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Compute the Difference LAPACK vs. Magma
=================================================================== */
error2 = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ylapack, &incx );
error2 = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work ) / error2;
}
/* =====================================================================
Compute the Difference Cublas vs. Magma
=================================================================== */
error = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ycublas, &incx );
error = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work ) / error;
bool okay = (error < tol && error2 < tol);
status += ! okay;
if ( opts.lapack ) {
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) N, (int) offset,
cpu_perf, cpu_time*1000.,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, error2, (okay ? "ok" : "failed") );
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e --- %s\n",
(int) N, (int) offset,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, (okay ? "ok" : "failed") );
}
/* Free Memory */
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_CPU( Ymagma1 );
TESTING_FREE_CPU( Ylapack );
TESTING_FREE_PIN( X );
TESTING_FREE_PIN( hwork );
magma_setdevice( opts.device );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
for( dev=0; dev < opts.ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
// comment out these two lines line & top of loop test a specific offset
} // end for ioffset
printf( "\n" );
}
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_destroy( queues[dev] );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, ldda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||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;
gflops = FLOPS_ZPOTRI( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* factorize matrix */
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_zpotrf_gpu( opts.uplo, N, d_A, ldda, &info );
// check for exact singularity
//magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
//h_R[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_R, lda, d_A, ldda );
gpu_time = magma_wtime();
magma_zpotri_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_zpotri( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
error = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / 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( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error;
magmaFloatComplex *h_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |Ax-b|/(N*|A|*|x|)\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |PA-LU|/(N*|A|)\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_CGETRF( M, N ) / 1e9;
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_PIN( h_A, magmaFloatComplex, n2 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_cgetrf(&M, &N, h_A, &lda, ipiv, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( M, N, h_A, lda );
gpu_time = magma_wtime();
magma_cgetrf( M, N, h_A, lda, ipiv, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time );
}
if ( opts.check == 2 ) {
error = get_residual( M, N, h_A, lda, ipiv );
printf(" %8.2e %s\n", error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else if ( opts.check ) {
error = get_LU_error( M, N, h_A, lda, ipiv );
printf(" %8.2e %s\n", error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_PIN( h_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, e1, e2, e3, e4, e5, *work;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, *h_rwork, tmp[1];
magmaFloatComplex_ptr 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;
opts.parse_opts( argc, argv );
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", (int) opts.version );
return -1;
}
float tol = 10. * opts.tolerance * lapackf77_slamch("E");
printf("%% version %d\n", (int) opts.version );
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 cgegqr requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because cgegqr requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
ldwork = N*N;
if (opts.version == 2) {
ldwork = 3*N*N + min_mn + 2;
}
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN(h_rwork, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( work, float, M );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, ldwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
if ( opts.warmup ) {
magma_cgegqr_gpu( 1, M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magmablasSetKernelStream( opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
magma_cgegqr_gpu( opts.version, M, N, d_A, ldda, dwork, h_rwork, &info );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_cgetmatrix( M, N, d_A, ldda, h_R, lda );
// Regenerate R
// blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_R, &lda, h_A, &lda, &c_zero, h_rwork, &N);
// magma_cprint(N, N, h_work, N);
blasf77_ctrmm("r", "u", "n", "n", &M, &N, &c_one, h_rwork, &N, h_R, &lda);
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
e5 = lapackf77_clange("i", &M, &N, h_R, &lda, work) /
lapackf77_clange("i", &M, &N, h_A, &lda, work);
magma_cgetmatrix( M, N, d_A, ldda, h_R, lda );
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&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_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_cgemm("c", "n", &N, &N, &M, &c_one, h_R, &lda, h_R, &lda, &c_zero, h_work, &N);
for (int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e3 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
blasf77_cgemm("c", "n", &N, &N, &M, &c_one, h_A, &lda, h_A, &lda, &c_zero, h_work, &N);
for (int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e4 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
if (opts.version != 4)
error = e1;
else
error = e1 / (10.*max(M,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,
(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_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" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time /*cpu_time*/;
magmaFloatComplex *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, nb, lwork, liwork;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
#ifdef COMPLEX
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_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");
// TODO: how to check NoVec? This doesn't have an equivalent call to LAPACK.
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "WARNING: cannot currently check results when not computing vectors (option -JN).\n" );
}
printf("using: itype = %d, jobz = %s, uplo = %s, check = %d, fraction = %6.4f\n",
(int) opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
(int) opts.check, opts.fraction);
printf(" N M GPU Time (sec)\n");
printf("============================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
nb = magma_get_chetrd_nb(N);
#ifdef COMPLEX
lwork = max( N + N*nb, 2*N + N*N );
lrwork = 1 + 5*N +2*N*N;
#else
lwork = max( 2*N + N*nb, 1 + 6*N + 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, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, 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, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_S, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
#ifdef COMPLEX
TESTING_MALLOC_PIN( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clarnv( &ione, ISEED, &n2, h_B );
magma_cmake_hpd( N, h_B, N );
magma_cmake_hermitian( N, h_A, N );
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if ( opts.warmup ) {
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_B, &N, h_S, &N );
magma_chegvdx( opts.itype, opts.jobz, MagmaRangeI, opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_chegvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_chegvdx( opts.itype, opts.jobz, MagmaRangeI, opts.uplo,
N, h_R, N, h_S, N, 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_chegvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check && opts.jobz != MagmaNoVec ) {
/* =====================================================================
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) - D(w/o V) | / | D |
=================================================================== */
#ifdef REAL
float *rwork = h_work + N*N;
#endif
result[0] = 1.;
result[0] /= lapackf77_clanhe("1", lapack_uplo_const(opts.uplo), &N, h_A, &N, rwork);
result[0] /= lapackf77_clange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_chemm("L", lapack_uplo_const(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_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_clange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_chemm("L", lapack_uplo_const(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_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_clange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_chemm("L", lapack_uplo_const(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_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_clange("1", &N, &m1, h_R, &N, rwork)/N;
}
// Disable eigenvalue check which calls routine again --
// it obscures whether error occurs in first call above or in this call.
// TODO: add comparison to LAPACK, as in testing_chegvd.cpp?
//
//lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &N, h_R, &N );
//lapackf77_clacpy( MagmaFullStr, &N, &N, h_B, &N, h_S, &N );
//
//magma_int_t m2;
//magma_chegvdx( opts.itype, MagmaNoVec, MagmaRangeI, opts.uplo,
// N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
// h_work, lwork,
// #ifdef COMPLEX
// rwork, lrwork,
// #endif
// iwork, liwork,
// &info );
//if (info != 0)
// printf("magma_chegvdx returned error %d: %s.\n",
// (int) info, magma_strerror( info ));
//
//float maxw=0, diff=0;
//for(int j=0; j < m2; j++) {
// maxw = max(maxw, fabs(w1[j]));
// maxw = max(maxw, fabs(w2[j]));
// diff = max(diff, fabs(w1[j]-w2[j]));
//}
//result[1] = diff / (m2*maxw);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %7.2f\n",
(int) N, (int) m1, gpu_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"));
}
//printf( " | D(w/ Z) - D(w/o Z) | / |D| = %8.2e %s\n\n", result[1], (result[1] < tolulp ? "ok" : "failed"));
status += ! (result[0] < tol); // && result[1] < tolulp);
}
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" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
magmaFloatComplex_ptr d_A;
magmaFloat_ptr 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_CGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, float, min_mn );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
if ( opts.warmup ) {
magma_cgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_cgeqr2_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_cgeqr2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&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_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
error = lapackf77_clange("f", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("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;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgesv_gpu
*/
int main(int argc , char **argv)
{
real_Double_t gflops, gpu_perf, gpu_time;
double Rnorm, Anorm, Xnorm, *work;
double *hA, *hB, *hX;
magmaDouble_ptr dA, dB;
magma_int_t *ipiv;
magma_int_t N = 0, n2, lda, ldb, ldda, lddb;
magma_int_t size[7] =
{ 1024, 2048, 3072, 4032, 5184, 6048, 7000};
magma_int_t i, info, szeB;
double z_one = MAGMA_D_ONE;
double mz_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t NRHS = 100;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
if (strcmp("-R", argv[i])==0)
NRHS = atoi(argv[++i]);
}
if (N>0) size[0] = size[6] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_dgesv_gpu -N <matrix size> -R <right hand sides>\n\n");
}
/* Initialize */
magma_queue_t queue;
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], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
/* Allocate memory for the largest matrix */
N = size[6];
n2 = N * N;
ldda = ((N+31)/32) * 32;
// ldda = N;
lddb = ldda;
TESTING_MALLOC_PIN( ipiv, magma_int_t, N );
TESTING_MALLOC_PIN( hA, double, n2 );
TESTING_MALLOC_PIN( hB, double, N*NRHS );
TESTING_MALLOC_PIN( hX, double, N*NRHS );
TESTING_MALLOC_PIN( work, double, N );
TESTING_MALLOC_DEV( dA, double, ldda*N );
TESTING_MALLOC_DEV( dB, double, lddb*NRHS );
printf("\n\n");
printf(" N NRHS GPU GFlop/s (sec) ||B - AX|| / ||A||*||X||\n");
printf("===========================================================\n");
for( i = 0; i < 7; i++ ) {
N = size[i];
lda = N;
ldb = lda;
n2 = lda*N;
szeB = ldb*NRHS;
ldda = ((N+31)/32)*32;
//ldda = N;
lddb = ldda;
gflops = ( FLOPS_GETRF( (double)N, (double)N ) +
FLOPS_GETRS( (double)N, (double)NRHS ) ) / 1e9;
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &n2, hA );
lapackf77_dlarnv( &ione, ISEED, &szeB, hB );
/* Warm up to measure the performance */
magma_dsetmatrix( N, N, hA, 0, lda, dA, 0, ldda, queue );
magma_dsetmatrix( N, NRHS, hB, 0, lda, dB, 0, lddb, queue );
magma_dgesv_gpu( N, NRHS, dA, 0, ldda, ipiv, dB, 0, lddb, &info, queue );
//=====================================================================
// Solve Ax = b through an LU factorization
//=====================================================================
magma_dsetmatrix( N, N, hA, 0, lda, dA, 0, ldda, queue );
magma_dsetmatrix( N, NRHS, hB, 0, lda, dB, 0, lddb, queue );
gpu_time = magma_wtime();
magma_dgesv_gpu( N, NRHS, dA, 0, ldda, ipiv, dB, 0, lddb, &info, queue );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf( "magma_dposv had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* =====================================================================
Residual
=================================================================== */
magma_dgetmatrix( N, NRHS, dB, 0, lddb, hX, 0, ldb, queue );
Anorm = lapackf77_dlange("I", &N, &N, hA, &lda, work);
Xnorm = lapackf77_dlange("I", &N, &NRHS, hX, &ldb, work);
blasf77_dgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &NRHS, &N,
&z_one, hA, &lda,
hX, &ldb,
&mz_one, hB, &ldb );
Rnorm = lapackf77_dlange("I", &N, &NRHS, hB, &ldb, work);
printf( "%5d %5d %7.2f (%7.2f) %8.2e\n",
N, NRHS, gpu_perf, gpu_time, Rnorm/(Anorm*Xnorm) );
if (argc != 1)
break;
}
/* clean up */
TESTING_FREE_PIN( hA );
TESTING_FREE_PIN( hB );
TESTING_FREE_PIN( hX );
TESTING_FREE_PIN( work );
TESTING_FREE_PIN( ipiv );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dB );
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cungqr
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float Anorm, error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *hA, *hR, *tau, *h_work;
magmaFloatComplex_ptr dA, dT;
magma_int_t m, n, k;
magma_int_t n2, lda, ldda, lwork, min_mn, nb, info;
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("Running version %d; available are (specified through --version num):\n",
(int) opts.version);
printf("1 - uses precomputed clarft matrices (default)\n");
printf("2 - recomputes the clarft matrices on the fly\n\n");
printf(" m n k CPU GFlop/s (sec) GPU GFlop/s (sec) ||R|| / ||A||\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];
k = opts.ksize[itest];
if ( m < n || n < k ) {
printf( "%5d %5d %5d skipping because m < n or n < k\n", (int) m, (int) n, (int) k );
continue;
}
lda = m;
ldda = ((m + 31)/32)*32;
n2 = lda*n;
min_mn = min(m, n);
nb = magma_get_cgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_CUNGQR( m, n, k ) / 1e9;
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN( hR, magmaFloatComplex, lda*n );
TESTING_MALLOC_CPU( hA, magmaFloatComplex, lda*n );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, ldda*n );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_clarnv( &ione, ISEED, &n2, hA );
lapackf77_clacpy( MagmaFullStr, &m, &n, hA, &lda, hR, &lda );
Anorm = lapackf77_clange("f", &m, &n, hA, &lda, work );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// first, get QR factors in both hA and hR
// okay that magma_cgeqrf_gpu has special structure for R; R isn't used here.
magma_csetmatrix( m, n, hA, lda, dA, ldda );
magma_cgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_cgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( m, n, dA, ldda, hA, lda );
lapackf77_clacpy( MagmaFullStr, &m, &n, hA, &lda, hR, &lda );
gpu_time = magma_wtime();
if (opts.version == 1)
magma_cungqr( m, n, k, hR, lda, tau, dT, nb, &info );
else
magma_cungqr2(m, n, k, hR, lda, tau, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cungqr( &m, &n, &k, hA, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_caxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_clange("f", &m, &n, hR, &lda, work) / Anorm;
bool okay = (error < tol);
status += ! okay;
printf("%5d %5d %5d %7.1f (%7.2f) %7.1f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (okay ? "ok" : "failed"));
}
else {
printf("%5d %5d %5d --- ( --- ) %7.1f (%7.2f) --- \n",
(int) m, (int) n, (int) k,
gpu_perf, gpu_time );
}
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( hR );
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zheevd_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
magmaDoubleComplex *h_A, *h_R, *d_R, *h_work, aux_work[1];
#ifdef COMPLEX
double *rwork, aux_rwork[1];
magma_int_t lrwork;
#endif
double *w1, *w2, result[4]={0, 0, 0, 0}, eps;
magma_int_t *iwork, aux_iwork[1];
magma_int_t N, n2, info, lwork, liwork, lda, ldda;
magma_int_t izero = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
eps = lapackf77_dlamch( "E" );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
// checking NoVec requires LAPACK
opts.lapack |= (opts.check && opts.jobz == MagmaNoVec);
printf("using: jobz = %s, uplo = %s\n",
lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo));
printf(" N CPU Time (sec) GPU Time (sec)\n");
printf("=======================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
lda = N;
ldda = roundup( N, opts.roundup ); // by default, round to multiple of 32
// query for workspace sizes
magma_zheevd_gpu( opts.jobz, opts.uplo,
N, NULL, ldda, NULL,
NULL, lda,
aux_work, -1,
#ifdef COMPLEX
aux_rwork, -1,
#endif
aux_iwork, -1,
&info );
lwork = (magma_int_t) MAGMA_Z_REAL( aux_work[0] );
#ifdef COMPLEX
lrwork = (magma_int_t) aux_rwork[0];
#endif
liwork = aux_iwork[0];
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, N*lda );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
#ifdef COMPLEX
TESTING_MALLOC_CPU( rwork, double, lrwork );
#endif
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, N*lda );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_DEV( d_R, magmaDoubleComplex, N*ldda );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hermitian( N, h_A, N );
magma_zsetmatrix( N, N, h_A, lda, d_R, ldda );
/* warm up run */
if ( opts.warmup ) {
magma_zheevd_gpu( opts.jobz, opts.uplo,
N, d_R, ldda, w1,
h_R, lda,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_zheevd_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_zsetmatrix( N, N, h_A, lda, d_R, ldda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zheevd_gpu( opts.jobz, opts.uplo,
N, d_R, ldda, w1,
h_R, lda,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_zheevd_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check && opts.jobz != MagmaNoVec ) {
/* =====================================================================
Check the results following the LAPACK's [zcds]drvst routine.
A is factored as A = U S U' and the following 3 tests computed:
(1) | A - U S U' | / ( |A| N )
(2) | I - U'U | / ( N )
(3) | S(with U) - S(w/o U) | / | S |
=================================================================== */
magma_zgetmatrix( N, N, d_R, ldda, h_R, lda );
magmaDoubleComplex *work;
TESTING_MALLOC_CPU( work, magmaDoubleComplex, 2*N*N );
// e=NULL is unused since kband=0; tau=NULL is unused since itype=1
lapackf77_zhet21( &ione, lapack_uplo_const(opts.uplo), &N, &izero,
h_A, &lda,
w1, NULL,
h_R, &lda,
h_R, &lda,
NULL, work,
#ifdef COMPLEX
rwork,
#endif
&result[0] );
result[0] *= eps;
result[1] *= eps;
TESTING_FREE_CPU( work ); work=NULL;
// Disable eigenvalue check which calls routine again --
// it obscures whether error occurs in first call above or in this call.
// But see comparison to LAPACK below.
//
//magma_zsetmatrix( N, N, h_A, lda, d_R, ldda );
//magma_zheevd_gpu( MagmaNoVec, opts.uplo,
// N, d_R, ldda, w2,
// h_R, lda,
// h_work, lwork,
// #ifdef COMPLEX
// rwork, lrwork,
// #endif
// iwork, liwork,
// &info);
//if (info != 0)
// printf("magma_zheevd_gpu returned error %d: %s.\n",
// (int) info, magma_strerror( info ));
//
//double 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[2] = diff / (N*maxw);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zheevd( lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
&N, h_A, &lda, w2,
h_work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork,
&info);
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_zheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compare eigenvalues
double 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[3] = diff / (N*maxw);
printf("%5d %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time);
}
else {
printf("%5d --- %7.2f\n",
(int) N, gpu_time);
}
/* =====================================================================
Print execution time
=================================================================== */
if ( opts.check && opts.jobz != MagmaNoVec ) {
printf("Testing the factorization A = U S U' for correctness:\n");
printf(" | A - U S U' | / (|A| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
printf(" | I - U'U | / N = %8.2e %s\n", result[1], (result[1] < tol ? "ok" : "failed") );
//printf(" | S(w/ U) - S(w/o U) | / |S| = %8.2e %s\n\n", result[2], (result[2] < tolulp ? "ok" : "failed") );
status += ! (result[0] < tol && result[1] < tol); // && result[2] < tolulp)
}
if ( opts.lapack ) {
printf(" | S_magma - S_lapack | / |S| = %8.2e %s\n\n", result[3], (result[3] < tolulp ? "ok" : "failed") );
status += ! (result[3] < tolulp);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
#ifdef COMPLEX
TESTING_FREE_CPU( rwork );
#endif
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( d_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time /*cpu_time*/;
magmaDoubleComplex *h_A, *h_R, *h_B, *h_S, *h_work;
double *w1, *w2, vl=0, vu=0;
double result[2] = {0};
magma_int_t *iwork;
magma_int_t N, n2, info, il, iu, m1, m2, nb, lwork, liwork;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
magma_int_t lrwork;
#endif
//double d_one = 1.;
//double d_ten = 10.;
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 );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: itype = %d, jobz = %s, uplo = %s, check = %d, fraction = %6.4f\n",
(int) opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
(int) opts.check, opts.fraction);
printf(" N M GPU Time (sec)\n");
printf("============================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
nb = magma_get_zhetrd_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, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_S, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, double, lrwork);
#endif
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
magma_zmake_hpd( N, h_B, N );
magma_zmake_hermitian( N, h_A, N );
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if(opts.warmup){
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_zhegvdx( opts.itype, opts.jobz, MagmaRangeI, 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_zhegvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_zhegvdx( opts.itype, opts.jobz, MagmaRangeI, 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_zhegvdx 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)
double *rwork = h_work + N*N;
#endif
double temp1, temp2;
result[0] = 1.;
result[0] /= lapackf77_zlanhe("1", lapack_uplo_const(opts.uplo), &N, h_A, &N, rwork);
result[0] /= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_zhemm("L", lapack_uplo_const(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_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_zhemm("L", lapack_uplo_const(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_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_zhemm("L", lapack_uplo_const(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_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_zhegvdx( opts.itype, MagmaNoVec, MagmaRangeI, 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_zhegvdx 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 / (((double)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 ? "ok" : "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 ? "ok" : "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 ? "ok" : "failed"));
}
printf( "(2) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e %s\n\n", result[1], (result[1] < tolulp ? "ok" : "failed"));
status += ! (result[0] < tol && result[1] < tolulp);
}
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
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zswap, zswapblk, zpermute, zlaswp, zlaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex *h_A1, *h_A2;
magmaDoubleComplex *d_A1, *d_A2;
magmaDoubleComplex *h_R1, *h_R2;
// row-major and column-major performance
real_Double_t row_perf0, col_perf0;
real_Double_t row_perf1, col_perf1;
real_Double_t row_perf2, col_perf2;
real_Double_t row_perf3;
real_Double_t row_perf4;
real_Double_t row_perf5, col_perf5;
real_Double_t row_perf6, col_perf6;
real_Double_t row_perf7;
real_Double_t cpu_perf;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magma_int_t *d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasZswap zswap zswapblk zlaswp zpermute zlaswp2 zlaswpx zcopymatrix CPU (all in )\n");
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj row-maj/col-maj row-blk/col-blk zlaswp (GByte/s)\n");
printf("==================================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_zgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaDoubleComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaDoubleComplex, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasZswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasZswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasZswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda);
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_zswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_zswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_zswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_zsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_zswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_zgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* zpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // zpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zpermute_long2( N, d_A1, ldda, ipiv2, nb, 0 );
time = magma_sync_wtime( queue ) - time;
row_perf3 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zlaswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_zlaswp2( N, d_A1, ldda, 1, nb, d_ipiv );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style zlaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zlaswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_zswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_zsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_zlaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_zlaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_zgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_zcopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_zcopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf3, ((check & 0x040) != 0 ? '*' : ' '),
row_perf4, ((check & 0x080) != 0 ? '*' : ' '),
row_perf7, ((check & 0x100) != 0 ? '*' : ' '),
row_perf5, ((check & 0x200) != 0 ? '*' : ' '),
col_perf5, ((check & 0x400) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf=0, cublas_time=0, cpu_perf, cpu_time;
double magma_error, cublas_error, magma_error2, cublas_error2;
magmaDoubleComplex *h_A, *h_R, *h_Amagma, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_A, *dtau_magma, *dtau_cublas;
magmaDoubleComplex **dA_array = NULL;
magmaDoubleComplex **dtau_array = NULL;
magma_int_t *dinfo_magma, *dinfo_cublas;
magma_int_t M, N, lda, ldda, lwork, n2, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_int_t batchCount;
magma_int_t column;
magma_opts opts( MagmaOptsBatched );
opts.parse_opts( argc, argv );
batchCount = opts.batchcount;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% BatchCount M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) |R - Q^H*A|_mag |I - Q^H*Q|_mag |R - Q^H*A|_cub |I - Q^H*Q|_cub\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 * batchCount;
ldda = M;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = (FLOPS_ZGEQRF( M, N ) + FLOPS_ZGEQRT( M, N )) / 1e9 * batchCount;
/* Allocate memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn * batchCount );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_Amagma, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( dtau_magma, magmaDoubleComplex, min_mn * batchCount);
TESTING_MALLOC_DEV( dtau_cublas, magmaDoubleComplex, min_mn * batchCount);
TESTING_MALLOC_DEV( dinfo_magma, magma_int_t, batchCount);
TESTING_MALLOC_DEV( dinfo_cublas, magma_int_t, batchCount);
TESTING_MALLOC_DEV( dA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dtau_array, magmaDoubleComplex*, batchCount );
// to determine the size of lwork
lwork = -1;
lapackf77_zgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_CPU( h_work, magmaDoubleComplex, lwork * batchCount);
column = N * batchCount;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaFullStr, &M, &column, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, column, h_R, lda, d_A, ldda );
magma_zset_pointer( dA_array, d_A, 1, 0, 0, ldda*N, batchCount, opts.queue );
magma_zset_pointer( dtau_array, dtau_magma, 1, 0, 0, min_mn, batchCount, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
info = magma_zgeqrf_batched(M, N, dA_array, ldda, dtau_array, dinfo_magma, batchCount, opts.queue);
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, column, d_A, ldda, h_Amagma, lda);
if (info != 0) {
printf("magma_zgeqrf_batched returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
/* cublasZgeqrfBatched is only available from CUBLAS v6.5 */
#if CUDA_VERSION >= 6050
magma_zsetmatrix( M, column, h_R, lda, d_A, ldda );
magma_zset_pointer( dA_array, d_A, 1, 0, 0, ldda*N, batchCount, opts.queue );
magma_zset_pointer( dtau_array, dtau_cublas, 1, 0, 0, min_mn, batchCount, opts.queue );
cublas_time = magma_sync_wtime( opts.queue );
int cublas_info; // not magma_int_t
cublasZgeqrfBatched( opts.handle, M, N, dA_array, ldda, dtau_array, &cublas_info, batchCount);
cublas_time = magma_sync_wtime( opts.queue ) - cublas_time;
cublas_perf = gflops / cublas_time;
if (cublas_info != 0) {
printf("cublasZgeqrfBatched returned error %d: %s.\n",
(int) cublas_info, magma_strerror( cublas_info ));
}
#endif
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
cpu_time = magma_wtime();
// #define BATCHED_DISABLE_PARCPU
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_int_t nthreads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
magma_set_omp_numthreads(nthreads);
#pragma omp parallel for schedule(dynamic)
#endif
for (magma_int_t s=0; s < batchCount; s++)
{
magma_int_t locinfo;
lapackf77_zgeqrf(&M, &N, h_A + s * lda * N, &lda, tau + s * min_mn, h_work + s * lwork, &lwork, &locinfo);
if (locinfo != 0) {
printf("lapackf77_zgeqrf matrix %d returned error %d: %s.\n",
(int) s, (int) locinfo, magma_strerror( locinfo ));
}
}
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_set_lapack_numthreads(nthreads);
#endif
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the MAGMA CUBLAS result compared to LAPACK
=================================================================== */
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
magmaDoubleComplex *Q, *R;
double *work;
TESTING_MALLOC_CPU( Q, magmaDoubleComplex, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, magmaDoubleComplex, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, double, min_mn );
/* check magma result */
magma_error = 0;
magma_error2 = 0;
magma_zgetvector(min_mn*batchCount, dtau_magma, 1, tau, 1);
for (int i=0; i < batchCount; i++)
{
double err, err2;
get_QR_error(M, N, min_mn,
h_Amagma + i*lda*N, h_R + i*lda*N, lda, tau + i*min_mn,
Q, ldq, R, ldr, h_work, lwork,
work, &err, &err2);
if ( isnan(err) || isinf(err) ) {
magma_error = err;
break;
}
magma_error = max( err, magma_error );
magma_error2 = max( err2, magma_error2 );
}
/* check cublas result */
cublas_error = 0;
cublas_error2 = 0;
#if CUDA_VERSION >= 6050
magma_zgetvector(min_mn*batchCount, dtau_magma, 1, tau, 1);
magma_zgetmatrix( M, column, d_A, ldda, h_A, lda);
for (int i=0; i < batchCount; i++)
{
double err, err2;
get_QR_error(M, N, min_mn,
h_A + i*lda*N, h_R + i*lda*N, lda, tau + i*min_mn,
Q, ldq, R, ldr, h_work, lwork,
work, &err, &err2);
if ( isnan(err) || isinf(err) ) {
cublas_error = err;
break;
}
cublas_error = max( err, cublas_error );
cublas_error2 = max( err2, cublas_error2 );
}
#endif
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
bool okay = (magma_error < tol && magma_error2 < tol);
//bool okay_cublas = (cublas_error < tol && cublas_error2 < tol);
status += ! okay;
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %15.2e %15.2e %15.2e %15.2e %s\n",
(int)batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, magma_error2,
cublas_error, cublas_error2,
(okay ? "ok" : "failed") );
}
else {
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) ---\n",
(int)batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_Amagma);
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau_magma );
TESTING_FREE_DEV( dtau_cublas );
TESTING_FREE_DEV( dinfo_magma );
TESTING_FREE_DEV( dinfo_cublas );
TESTING_FREE_DEV( dA_array );
TESTING_FREE_DEV( dtau_array );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A, *h_R;
magma_int_t N, n2, lda, info;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu = %d, uplo = %s\n", (int) opts.ngpu, lapack_uplo_const(opts.uplo) );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma - R_lapack||_F / ||R_lapack||_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;
gflops = FLOPS_CPOTRF( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
magma_cmake_hpd( N, h_A, lda );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cpotrf( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_clange("f", &N, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &N, &N, h_R, &lda, work) / 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( h_A );
TESTING_FREE_PIN( h_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhetrd_he2hb
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_time, gpu_perf;
magmaDoubleComplex *h_A, *h_R, *h_work;
magmaDoubleComplex *tau;
double *D, *E;
magma_int_t N, n2, lda, ldda, lwork, ldt, info, nstream;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
// TODO add these options to parse_opts
magma_int_t NE = 0;
magma_int_t distblk = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t WANTZ = (opts.jobz == MagmaVec);
double tol = opts.tolerance * lapackf77_dlamch("E");
if (opts.nb == 0)
opts.nb = 64; //magma_get_zhetrd_he2hb_nb(N);
if (NE < 1)
NE = N; //64; //magma_get_zhetrd_he2hb_nb(N);
nstream = max(3, opts.ngpu+2);
magma_queue_t streams[MagmaMaxGPUs][20];
magmaDoubleComplex_ptr da[MagmaMaxGPUs], dT1[MagmaMaxGPUs];
if ((distblk == 0) || (distblk < opts.nb))
distblk = max(256, opts.nb);
printf("voici ngpu %d distblk %d NB %d nstream %d\n ",
(int) opts.ngpu, (int) distblk, (int) opts.nb, (int) nstream);
for( magma_int_t dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( int i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[dev][i] );
}
}
magma_setdevice( 0 );
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldt = N;
ldda = ((N+31)/32)*32;
n2 = N*lda;
/* We suppose the magma NB is bigger than lapack NB */
lwork = N*opts.nb;
//gflops = ....?
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, N-1 );
TESTING_MALLOC_PIN( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_PIN( D, double, N );
TESTING_MALLOC_PIN( E, double, N );
for( magma_int_t dev = 0; dev < opts.ngpu; ++dev ) {
magma_int_t mlocal = ((N / distblk) / opts.ngpu + 1) * distblk;
magma_setdevice( dev );
TESTING_MALLOC_DEV( da[dev], magmaDoubleComplex, ldda*mlocal );
TESTING_MALLOC_DEV( dT1[dev], magmaDoubleComplex, N*opts.nb );
}
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hermitian( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* Copy the matrix to the GPU */
magma_zsetmatrix_1D_col_bcyclic( N, N, h_R, lda, da, ldda, opts.ngpu, distblk);
//magmaDoubleComplex_ptr dabis;
//TESTING_MALLOC_DEV( dabis, magmaDoubleComplex, ldda*N );
//magma_zsetmatrix(N, N, h_R, lda, dabis, ldda);
for (int count=0; count < 1; ++count) {
magma_setdevice(0);
gpu_time = magma_wtime();
if (opts.version == 30) {
magma_zhetrd_he2hb_mgpu_spec(
opts.uplo, N, opts.nb, h_R, lda, tau, h_work, lwork,
da, ldda, dT1, opts.nb, opts.ngpu, distblk,
streams, nstream, opts.nthread, &info);
} else {
nstream = 3;
magma_zhetrd_he2hb_mgpu(
opts.uplo, N, opts.nb, h_R, lda, tau, h_work, lwork,
da, ldda, dT1, opts.nb, opts.ngpu, distblk,
streams, nstream, opts.nthread, &info);
}
// magma_zhetrd_he2hb(opts.uplo, N, opts.nb, h_R, lda, tau, h_work, lwork, dT1[0], &info);
gpu_time = magma_wtime() - gpu_time;
printf(" Finish BAND N %d NB %d dist %d ngpu %d version %d timing= %f\n",
N, opts.nb, distblk, opts.ngpu, opts.version, gpu_time);
}
magma_setdevice(0);
for( magma_int_t dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice(dev);
magma_device_sync();
}
magma_setdevice(0);
magmablasSetKernelStream( NULL );
// todo neither of these is declared in headers
// magma_zhetrd_bhe2trc_v5(opts.nthread, WANTZ, opts.uplo, NE, N, opts.nb, h_R, lda, D, E, dT1[0], ldt);
// magma_zhetrd_bhe2trc(opts.nthread, WANTZ, opts.uplo, NE, N, opts.nb, h_R, lda, D, E, dT1[0], ldt);
// todo where is this timer started?
// gpu_time = magma_wtime() - gpu_time;
// todo what are the gflops?
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zhetrd_he2hb returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Print performance and error.
=================================================================== */
#if defined(CHECKEIG)
#if defined(PRECISION_z) || defined(PRECISION_d)
if ( opts.check ) {
printf(" Total N %5d flops %6.2f timing %6.2f seconds\n", (int) N, gpu_perf, gpu_time );
char JOBZ;
if (WANTZ == 0)
JOBZ = 'N';
else
JOBZ = 'V';
double nrmI=0.0, nrm1=0.0, nrm2=0.0;
int lwork2 = 256*N;
magmaDoubleComplex *work2, *AINIT;
double *rwork2, *D2;
// TODO free this memory !
magma_zmalloc_cpu( &work2, lwork2 );
magma_dmalloc_cpu( &rwork2, N );
magma_dmalloc_cpu( &D2, N );
magma_zmalloc_cpu( &AINIT, N*lda );
memcpy(AINIT, h_A, N*lda*sizeof(magmaDoubleComplex));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
int nt = min(12, opts.nthread);
#if defined(USEMKL)
mkl_set_num_threads(nt);
#endif
#if defined(USEACML)
omp_set_num_threads(nt);
#endif
#if defined(PRECISION_z) || defined (PRECISION_c)
lapackf77_zheev( "N", "L", &N, h_A, &lda, D2, work2, &lwork2, rwork2, &info );
#else
lapackf77_dsyev( "N", "L", &N, h_A, &lda, D2, work2, &lwork2, &info );
#endif
///* call eigensolver for our resulting tridiag [D E] and for Q */
//dstedc_withZ('V', N, D, E, h_R, lda);
////dsterf_( &N, D, E, &info);
cpu_time = magma_wtime() - cpu_time;
printf(" Finish CHECK - EIGEN timing= %f threads %d\n", cpu_time, nt);
/* compare result */
cmp_vals(N, D2, D, &nrmI, &nrm1, &nrm2);
magmaDoubleComplex *WORKAJETER;
double *RWORKAJETER, *RESU;
// TODO free this memory !
magma_zmalloc_cpu( &WORKAJETER, (2* N * N + N) );
magma_dmalloc_cpu( &RWORKAJETER, N );
magma_dmalloc_cpu( &RESU, 10 );
int MATYPE;
memset(RESU, 0, 10*sizeof(double));
MATYPE=3;
double NOTHING=0.0;
cpu_time = magma_wtime();
// check results
zcheck_eig_( lapack_vec_const(opts.jobz), &MATYPE, &N, &opts.nb,
AINIT, &lda, &NOTHING, &NOTHING, D2, D,
h_R, &lda, WORKAJETER, RWORKAJETER, RESU );
cpu_time = magma_wtime() - cpu_time;
printf(" Finish CHECK - results timing= %f\n", cpu_time);
#if defined(USEMKL)
mkl_set_num_threads(1);
#endif
#if defined(USEACML)
omp_set_num_threads(1);
#endif
printf("\n");
printf(" ================================================================================================================\n");
printf(" ==> INFO voici threads=%d N=%d NB=%d WANTZ=%d\n", (int) opts.nthread, (int) N, (int) opts.nb, (int) WANTZ);
printf(" ================================================================================================================\n");
printf(" DSBTRD : %15s \n", "STATblgv9withQ ");
printf(" ================================================================================================================\n");
if (WANTZ > 0)
printf(" | A - U S U' | / ( |A| n ulp ) : %15.3E \n", RESU[0]);
if (WANTZ > 0)
printf(" | I - U U' | / ( n ulp ) : %15.3E \n", RESU[1]);
printf(" | D1 - EVEIGS | / (|D| ulp) : %15.3E \n", RESU[2]);
printf(" max | D1 - EVEIGS | : %15.3E \n", RESU[6]);
printf(" ================================================================================================================\n\n\n");
printf(" ****************************************************************************************************************\n");
printf(" * Hello here are the norm Infinite (max)=%8.2e norm one (sum)=%8.2e norm2(sqrt)=%8.2e *\n", nrmI, nrm1, nrm2);
printf(" ****************************************************************************************************************\n\n");
}
#endif // PRECISION_z || PRECISION_d
#endif // CHECKEIG
printf(" Total N %5d flops %6.2f timing %6.2f seconds\n", (int) N, 0.0, gpu_time );
printf("============================================================================\n\n\n");
TESTING_FREE_CPU( tau );
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( D );
TESTING_FREE_PIN( E );
for( magma_int_t dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
TESTING_FREE_DEV( da[dev] );
TESTING_FREE_DEV( dT1[dev] );
}
magma_setdevice( 0 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
for( magma_int_t dev = 0; dev < opts.ngpu; ++dev ) {
for( int i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[dev][i] );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e1, e2, work[1];
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
magmaFloatComplex *d_A, *dwork, *ddA, *d_T;
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_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I - Q'Q||_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;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( d_T, magmaFloatComplex, N*N );
cudaMemset( ddA, 0, N*N*sizeof(magmaFloatComplex) );
cudaMemset( d_T, 0, N*N*sizeof(magmaFloatComplex) );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgegqr_gpu( M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
if (opts.version == 2) {
int min_mn = min(M, N);
int nb = N;
cuFloatComplex *dtau = dwork;
magma_cgeqr2x3_gpu(&M, &N, d_A, &ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn), &info);
magma_cgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_cungqr_gpu( M, N, N, d_A, ldda, tau, d_T, nb, &info );
}
else
magma_cgegqr_gpu( M, N, d_A, ldda, dwork, h_work, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&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_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
magmaFloatComplex one = MAGMA_C_ONE, zero = MAGMA_C_ZERO;
blasf77_cgemm("t", "n", &N, &N, &M, &one, h_R, &M, h_R, &M, &zero, h_work, &N);
for(int ii=0; ii<N*N; ii+=(N+1)) h_work[ii] = MAGMA_C_SUB(h_work[ii], one);
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work);
blasf77_cgemm("t", "n", &N, &N, &M, &one, h_A, &M, h_A, &M, &zero, h_work, &N);
for(int ii=0; ii<N*N; ii+=(N+1)) h_work[ii] = MAGMA_C_SUB(h_work[ii], one);
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, e1, e2 );
}
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_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( d_T );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magma_int_t N, n2, lda, ldda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magmaDoubleComplex **d_A_array = NULL;
magma_int_t *dinfo_magma;
magma_int_t batchCount;
magma_queue_t queue = magma_stream;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
batchCount = opts.batchcount;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("BatchCount N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
ldda = lda = ((N+31)/32)*32;
n2 = lda* N * batchCount;
gflops = batchCount * FLOPS_ZPOTRF( N ) / 1e9 ;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2);
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2);
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda * N * batchCount);
TESTING_MALLOC_DEV( dinfo_magma, magma_int_t, batchCount);
magma_malloc((void**)&d_A_array, batchCount * sizeof(*d_A_array));
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
for(int i=0; i<batchCount; i++)
{
magma_zmake_hpd( N, h_A + i * lda * N, lda );// need modification
}
magma_int_t columns = N * batchCount;
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &(columns), h_A, &lda, h_R, &lda );
magma_zsetmatrix( N, columns, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
zset_pointer(d_A_array, d_A, ldda, 0, 0, ldda * N, batchCount, queue);
gpu_time = magma_sync_wtime(NULL);
info = magma_zpotrf_batched( opts.uplo, N, d_A_array, ldda, dinfo_magma, batchCount, queue);
gpu_time = magma_sync_wtime(NULL) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_int_t *cpu_info = (magma_int_t*) malloc(batchCount*sizeof(magma_int_t));
magma_getvector( batchCount, sizeof(magma_int_t), dinfo_magma, 1, cpu_info, 1);
for(int i=0; i<batchCount; i++)
{
if(cpu_info[i] != 0 ){
printf("magma_zpotrf_batched matrix %d returned internal error %d\n",i, (int)cpu_info[i] );
}
}
if (info != 0)
printf("magma_zpotrf_batched returned argument error %d: %s.\n", (int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for(int i=0; i<batchCount; i++)
{
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A + i * lda * N, &lda, &info );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, columns, d_A, ldda, h_R, lda );
magma_int_t NN = lda*N;
char const uplo = 'l'; // lapack_uplo_const(opts.uplo)
double err = 0.0;
for(int i=0; i<batchCount; i++)
{
error = lapackf77_zlanhe("f", &uplo, &N, h_A + i * lda*N, &lda, work);
blasf77_zaxpy(&NN, &c_neg_one, h_A + i * lda*N, &ione, h_R + i * lda*N, &ione);
error = lapackf77_zlanhe("f", &uplo, &N, h_R + i * lda*N, &lda, work) / error;
if ( isnan(error) || isinf(error) ) {
err = error;
break;
}
err = max(fabs(error),err);
}
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int)batchCount, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., err, (error < tol ? "ok" : "failed"));
status += ! (err < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int)batchCount, (int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_A_array );
TESTING_FREE_DEV( dinfo_magma );
free(cpu_info);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cheevd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
/* Constants */
const float d_zero = 0;
const magma_int_t izero = 0;
const magma_int_t ione = 1;
/* Local variables */
real_Double_t gpu_time, cpu_time;
magmaFloatComplex *h_A, *h_R, *h_Z, *h_work, aux_work[1];
#ifdef COMPLEX
float *rwork, aux_rwork[1];
magma_int_t lrwork;
#endif
float *w1, *w2, result[4]={0, 0, 0, 0}, eps, abstol;
magma_int_t *iwork, *isuppz, *ifail, aux_iwork[1];
magma_int_t N, n2, info, lwork, liwork, lda;
magma_int_t ISEED[4] = {0,0,0,1};
eps = lapackf77_slamch( "E" );
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
// checking NoVec requires LAPACK
opts.lapack |= (opts.check && opts.jobz == MagmaNoVec);
magma_range_t range = MagmaRangeAll;
if (opts.fraction != 1)
range = MagmaRangeI;
#ifdef REAL
if (opts.version == 3 || opts.version == 4) {
printf("%% magma_cheevr and magma_cheevx are not available for real precisions (single, float).\n");
return status;
}
#endif
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
// pass ngpu = -1 to test multi-GPU code using 1 gpu
magma_int_t abs_ngpu = abs( opts.ngpu );
printf("%% jobz = %s, range = %s, uplo = %s, fraction = %6.4f, ngpu = %d\n",
lapack_vec_const(opts.jobz), lapack_range_const(range), lapack_uplo_const(opts.uplo),
opts.fraction, int(abs_ngpu) );
printf("%% N CPU Time (sec) GPU Time (sec) |S-S_magma| |A-USU^H| |I-U^H U|\n");
printf("%%============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
lda = N;
abstol = 0; // auto, in cheevr
// 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) );
}
// query for workspace sizes
if ( opts.version == 1 || opts.version == 2 ) {
magma_cheevd( opts.jobz, opts.uplo,
N, NULL, lda, NULL, // A, w
aux_work, -1,
#ifdef COMPLEX
aux_rwork, -1,
#endif
aux_iwork, -1,
&info );
}
else if ( opts.version == 3 ) {
#ifdef COMPLEX
magma_cheevr( opts.jobz, range, opts.uplo,
N, NULL, lda, // A
vl, vu, il, iu, abstol,
&m1, NULL, // w
NULL, lda, NULL, // Z, isuppz
aux_work, -1,
#ifdef COMPLEX
aux_rwork, -1,
#endif
aux_iwork, -1,
&info );
#endif
}
else if ( opts.version == 4 ) {
#ifdef COMPLEX
magma_cheevx( opts.jobz, range, opts.uplo,
N, NULL, lda, // A
vl, vu, il, iu, abstol,
&m1, NULL, // w
NULL, lda, // Z
aux_work, -1,
#ifdef COMPLEX
aux_rwork,
#endif
aux_iwork,
NULL, // ifail
&info );
// cheevx doesn't query rwork, iwork; set them for consistency
aux_rwork[0] = float(7*N);
aux_iwork[0] = float(5*N);
#endif
}
lwork = (magma_int_t) MAGMA_C_REAL( aux_work[0] );
#ifdef COMPLEX
lrwork = (magma_int_t) aux_rwork[0];
#endif
liwork = aux_iwork[0];
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, N*lda );
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
#ifdef COMPLEX
TESTING_MALLOC_CPU( rwork, float, lrwork );
#endif
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, N*lda );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
if (opts.version == 3) {
TESTING_MALLOC_CPU( h_Z, magmaFloatComplex, N*lda );
TESTING_MALLOC_CPU( isuppz, magma_int_t, 2*max(1,N) );
}
if (opts.version == 4) {
TESTING_MALLOC_CPU( h_Z, magmaFloatComplex, N*lda );
TESTING_MALLOC_CPU( ifail, magma_int_t, N );
}
/* Clear eigenvalues, for |S-S_magma| check when fraction < 1. */
lapackf77_slaset( "Full", &N, &ione, &d_zero, &d_zero, w1, &N );
lapackf77_slaset( "Full", &N, &ione, &d_zero, &d_zero, w2, &N );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
magma_cmake_hermitian( N, h_A, lda );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
if (opts.version == 1) {
if (opts.ngpu == 1) {
magma_cheevd( opts.jobz, opts.uplo,
N, h_R, lda, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
else {
//printf( "magma_cheevd_m, ngpu %d (%d)\n", opts.ngpu, abs_ngpu );
magma_cheevd_m( abs_ngpu, opts.jobz, opts.uplo,
N, h_R, lda, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
}
else if ( opts.version == 2 ) { // version 2: cheevdx computes selected eigenvalues/vectors
if (opts.ngpu == 1) {
magma_cheevdx( opts.jobz, range, opts.uplo,
N, h_R, lda,
vl, vu, il, iu,
&m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
else {
//printf( "magma_cheevdx_m, ngpu %d (%d)\n", opts.ngpu, abs_ngpu );
magma_cheevdx_m( abs_ngpu, opts.jobz, range, opts.uplo,
N, h_R, lda,
vl, vu, il, iu,
&m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
//printf( "il %d, iu %d, m1 %d\n", il, iu, m1 );
}
else if ( opts.version == 3 ) { // version 3: MRRR, computes selected eigenvalues/vectors
// only complex version available
#ifdef COMPLEX
magma_cheevr( opts.jobz, range, opts.uplo,
N, h_R, lda,
vl, vu, il, iu, abstol,
&m1, w1,
h_Z, lda, isuppz,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
lapackf77_clacpy( "Full", &N, &N, h_Z, &lda, h_R, &lda );
#endif
}
else if ( opts.version == 4 ) { // version 3: cheevx (QR iteration), computes selected eigenvalues/vectors
// only complex version available
#ifdef COMPLEX
magma_cheevx( opts.jobz, range, opts.uplo,
N, h_R, lda,
vl, vu, il, iu, abstol,
&m1, w1,
h_Z, lda,
h_work, lwork,
#ifdef COMPLEX
rwork, /*lrwork,*/
#endif
iwork, /*liwork,*/
ifail,
&info );
lapackf77_clacpy( "Full", &N, &N, h_Z, &lda, h_R, &lda );
#endif
}
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_cheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
bool okay = true;
if ( opts.check && opts.jobz != MagmaNoVec ) {
/* =====================================================================
Check the results following the LAPACK's [zcds]drvst routine.
A is factored as A = U S U^H and the following 3 tests computed:
(1) | A - U S U^H | / ( |A| N )
(2) | I - U^H U | / ( N )
(3) | S(with U) - S(w/o U) | / | S | // currently disabled, but compares to LAPACK
=================================================================== */
magmaFloatComplex *work;
TESTING_MALLOC_CPU( work, magmaFloatComplex, 2*N*N );
// e=NULL is unused since kband=0; tau=NULL is unused since itype=1
lapackf77_chet21( &ione, lapack_uplo_const(opts.uplo), &N, &izero,
h_A, &lda,
w1, NULL,
h_R, &lda,
h_R, &lda,
NULL, work,
#ifdef COMPLEX
rwork,
#endif
&result[0] );
result[0] *= eps;
result[1] *= eps;
TESTING_FREE_CPU( work ); work=NULL;
// Disable third eigenvalue check that calls routine again --
// it obscures whether error occurs in first call above or in this call.
// But see comparison to LAPACK below.
//
//lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
//magma_cheevd( MagmaNoVec, opts.uplo,
// N, h_R, lda, w2,
// h_work, lwork,
// #ifdef COMPLEX
// rwork, lrwork,
// #endif
// iwork, liwork,
// &info );
//if (info != 0) {
// printf("magma_cheevd 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[2] = diff / (N*maxw);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
if ( opts.version == 1 || opts.version == 2 ) {
lapackf77_cheevd( lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
&N, h_A, &lda, w2,
h_work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork,
&info );
}
else if ( opts.version == 3 ) {
lapackf77_cheevr( lapack_vec_const(opts.jobz),
lapack_range_const(range),
lapack_uplo_const(opts.uplo),
&N, h_A, &lda,
&vl, &vu, &il, &iu, &abstol,
&m1, w2,
h_Z, &lda, isuppz,
h_work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork,
&info );
lapackf77_clacpy( "Full", &N, &N, h_Z, &lda, h_A, &lda );
}
else if ( opts.version == 4 ) {
lapackf77_cheevx( lapack_vec_const(opts.jobz),
lapack_range_const(range),
lapack_uplo_const(opts.uplo),
&N, h_A, &lda,
&vl, &vu, &il, &iu, &abstol,
&m1, w2,
h_Z, &lda,
h_work, &lwork,
#ifdef COMPLEX
rwork,
#endif
iwork,
ifail,
&info );
lapackf77_clacpy( "Full", &N, &N, h_Z, &lda, h_A, &lda );
}
cpu_time = magma_wtime() - cpu_time;
if (info != 0) {
printf("lapackf77_cheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
// compare eigenvalues
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[3] = diff / (N*maxw);
okay = okay && (result[3] < tolulp);
printf("%5d %9.4f %9.4f %8.2e ",
(int) N, cpu_time, gpu_time, result[3] );
}
else {
printf("%5d --- %9.4f --- ",
(int) N, gpu_time);
}
// print error checks
if ( opts.check && opts.jobz != MagmaNoVec ) {
okay = okay && (result[0] < tol) && (result[1] < tol);
printf(" %8.2e %8.2e", result[0], result[1] );
}
else {
printf(" --- --- ");
}
printf(" %s\n", (okay ? "ok" : "failed"));
status += ! okay;
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
#ifdef COMPLEX
TESTING_FREE_CPU( rwork );
#endif
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
if ( opts.version == 3 ) {
TESTING_FREE_CPU( h_Z );
TESTING_FREE_CPU( isuppz );
}
if ( opts.version == 4 ) {
TESTING_FREE_CPU( h_Z );
TESTING_FREE_CPU( ifail );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_T, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex *d_A, *d_T, *ddA, *dtau;
magmaFloatComplex *d_A2, *d_T2, *ddA2, *dtau2;
float *dwork, *dwork2;
magma_int_t M, N, lda, ldda, lwork, n2, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
#define BLOCK_SIZE 64
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = 10. * opts.tolerance * lapackf77_slamch("E");
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
printf("version %d\n", (int) opts.version );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F/||A||_F ||R_T||\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 cgeqr2x requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because cgeqr2x 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_CGEQRF( M, N ) + FLOPS_CGEQRT( M, N )) / 1e9;
/* Allocate memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_T, magmaFloatComplex, N*N );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_T, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_T2, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA2, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( dtau2, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, float, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
TESTING_MALLOC_DEV( dwork2, float, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
// todo replace with magma_claset
cudaMemset(ddA, 0, N*N*sizeof(magmaFloatComplex));
cudaMemset(d_T, 0, N*N*sizeof(magmaFloatComplex));
cudaMemset(ddA2, 0, N*N*sizeof(magmaFloatComplex));
cudaMemset(d_T2, 0, N*N*sizeof(magmaFloatComplex));
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
magma_csetmatrix( M, N, h_R, lda, d_A2, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime(0);
if (opts.version == 1)
magma_cgeqr2x_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 2)
magma_cgeqr2x2_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 3)
magma_cgeqr2x3_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else {
printf( "call magma_cgeqr2x4_gpu\n" );
/*
Going through NULL stream is faster
Going through any stream is slower
Doing two streams in parallel is slower than doing them sequentially
Queuing happens on the NULL stream - user defined buffers are smaller?
*/
magma_cgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, NULL);
//magma_cgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, stream[1]);
//magma_cgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, stream[0]);
//magma_cgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, NULL);
//gflops *= 2;
}
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cgeqr2x_gpu version %d returned error %d: %s.\n",
(int) opts.version, (int) info, magma_strerror( info ));
}
else {
if ( opts.check ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&M, &N, h_A, &lda, tau, h_work, &N);
//magma_cgeqr2(&M, &N, h_A, &lda, tau, h_work, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
magma_cgetmatrix( N, N, ddA, N, h_T, N );
// Restore the upper triangular part of A before the check
for(int col=0; col < N; col++){
for(int row=0; row <= col; row++)
h_R[row + col*M] = h_T[row + col*N];
}
error = lapackf77_clange("M", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("M", &M, &N, h_R, &lda, work) / (N * error);
// Check if T is the same
magma_cgetmatrix( N, N, d_T, N, h_T, N );
float terr = 0.;
for(int col=0; col < N; col++)
for(int row=0; row <= col; row++)
terr += ( MAGMA_C_ABS(h_work[row + col*N] - h_T[row + col*N])*
MAGMA_C_ABS(h_work[row + col*N] - h_T[row + col*N]) );
terr = magma_ssqrt(terr);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, terr, (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_T );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_T );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( d_A2 );
TESTING_FREE_DEV( d_T2 );
TESTING_FREE_DEV( ddA2 );
TESTING_FREE_DEV( dtau2 );
TESTING_FREE_DEV( dwork2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegvd
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
magmaFloatComplex *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;
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;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_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");
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, 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_chetrd_nb(N);
#if defined(PRECISION_z) || defined(PRECISION_c)
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, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_B, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, 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 );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clarnv( &ione, ISEED, &n2, h_B );
magma_cmake_hpd( N, h_B, N );
magma_cmake_hermitian( N, h_A, N );
if ( opts.warmup || opts.check ) {
TESTING_MALLOC_CPU( h_Ainit, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_Binit, magmaFloatComplex, n2 );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_Binit, &N );
}
if (opts.warmup) {
// ==================================================================
// Warmup using MAGMA.
// ==================================================================
magma_chegvd_m( opts.ngpu, opts.itype, opts.jobz, opts.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_clacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
mgpu_time = magma_wtime();
magma_chegvd_m( opts.ngpu, opts.itype, opts.jobz, opts.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);
mgpu_time = magma_wtime() - mgpu_time;
if (info != 0)
printf("magma_chegvd_m returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
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_clanhe("1", lapack_uplo_const(opts.uplo), &N, h_Ainit, &N, rwork);
result /= lapackf77_clange("1", &N, &N, h_A, &N, rwork);
if (opts.itype == 1) {
blasf77_chemm("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_csscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result *= lapackf77_clange("1", &N, &N, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_chemm("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_csscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_clange("1", &N, &N, h_A, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_chemm("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_csscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_clange("1", &N, &N, h_A, &N, rwork)/N;
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_chegvd(opts.itype, opts.jobz, opts.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);
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_chegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_chegvd(&opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.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);
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_chegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
float temp1 = 0;
float temp2 = 0;
for(int j=0; j<N; j++) {
temp1 = max(temp1, fabs(w1[j]));
temp1 = max(temp1, fabs(w2[j]));
temp2 = max(temp2, fabs(w1[j]-w2[j]));
}
float result2 = temp2 / (((float)N)*temp1);
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %7.2f %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time, mgpu_time);
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( "(3) | D(MGPU)-D(LAPACK) |/ |D| = %8.2e %s\n\n", result2, (result2 < tolulp ? "ok" : "failed") );
status += ! (result < tol && result2 < tolulp);
}
else {
printf("%5d --- --- %7.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 || 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)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
magmaDoubleComplex *hA, *hR;
magmaDoubleComplex_ptr dA;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] =
{ 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 8160, 8192 };
magma_int_t i, info;
magmaDoubleComplex mz_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], matnorm, diffnorm;
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;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_zpotrf2_gpu -N %d\n\n", 1024);
}
/* Initialize */
magma_queue_t queue1, queue2;
magma_device_t device;
magma_int_t num = 0;
magma_int_t err;
magma_init();
err = magma_getdevices( &device, 2, &num );
if ( err != 0 or num < 1 ) {
fprintf( stderr, "magma_getdevices failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( device, &queue1 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( device, &queue2 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
magma_queue_t queues[2] = {queue1, queue2};
/* Allocate memory for the largest matrix */
N = size[9];
n2 = N * N;
ldda = ((N+31)/32) * 32;
TESTING_MALLOC_CPU( hA, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( hR, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*N );
printf("\n\n");
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( (double)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, hA );
/* Symmetrize and increase the diagonal */
for( int i = 0; i < N; ++i ) {
hA(i,i) = MAGMA_Z_MAKE( MAGMA_Z_REAL(hA(i,i)) + N, 0 );
for( int j = 0; j < i; ++j ) {
hA(i, j) = MAGMA_Z_CNJG( hA(j,i) );
}
}
lapackf77_zlacpy( MagmaFullStr, &N, &N, hA, &lda, hR, &lda );
/* Warm up to measure the performance */
magma_zsetmatrix( N, N, hA, lda, dA, 0, ldda, queue1);
clFinish(queue1);
magma_zpotrf2_gpu( MagmaLower, N, dA, 0, ldda, queues, &info );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( N, N, hA, lda, dA, 0, ldda, queue1 );
clFinish(queue1);
gpu_time = magma_wtime();
magma_zpotrf2_gpu( MagmaLower, N, dA, 0, ldda, queues, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf( "magma_zpotrf2 had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( MagmaLowerStr, &N, hA, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf( "lapackf77_zpotrf had error %d.\n", info );
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result compared to LAPACK
|R_magma - R_lapack| / |R_lapack|
=================================================================== */
magma_zgetmatrix( N, N, dA, 0, ldda, hR, lda, queue1 );
matnorm = lapackf77_zlange("f", &N, &N, hA, &lda, work);
blasf77_zaxpy(&n2, &mz_one, hA, &ione, hR, &ione);
diffnorm = lapackf77_zlange("f", &N, &N, hR, &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 (argc != 1)
break;
}
/* clean up */
TESTING_FREE_CPU( hA );
TESTING_FREE_PIN( hR );
TESTING_FREE_DEV( dA );
magma_queue_destroy( queue1 );
magma_queue_destroy( queue2 );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotf2_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex_ptr d_A;
magma_int_t N, n2, lda, ldda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double Anorm, error, work[1];
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% N CPU Gflop/s (ms) GPU Gflop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_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 = magma_roundup( N, opts.align ); // multiple of 32 by default
gflops = FLOPS_ZPOTRF( N ) / 1e9;
if ( N > 512 ) {
printf( "%5d skipping because zpotf2 does not support N > 512\n", (int) N );
continue;
}
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( opts.queue );
magma_zpotf2_gpu( opts.uplo, N, d_A, ldda, opts.queue, &info );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_zpotf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda, opts.queue );
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
Anorm = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / Anorm;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chesv_nopiv_gpu
*/
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_perf, cpu_time, gpu_perf, gpu_time;
float error, Rnorm, Anorm, Xnorm, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_B, *h_X, temp, *hwork;
magmaFloatComplex_ptr d_A, d_B;
magma_int_t *ipiv;
magma_int_t N, nrhs, lda, ldb, ldda, lddb, info, sizeA, sizeB, lwork;
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");
nrhs = opts.nrhs;
printf(" N NRHS CPU GFlop/s (sec) GPU GFlop/s (sec) ||B - AX|| / N*||A||*||X||\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;
ldb = lda;
ldda = ((N+31)/32)*32;
lddb = ldda;
gflops = ( FLOPS_CGETRF( N, N ) + FLOPS_CGETRS( N, nrhs ) ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( work, float, N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*nrhs );
/* Initialize the matrices */
sizeA = lda*N;
sizeB = ldb*nrhs;
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
bool nopiv = true;
if ( nopiv ) {
magma_cmake_hpd( N, h_A, lda ); // SPD / HPD does not require pivoting
}
else {
magma_cmake_hermitian( N, h_A, lda ); // symmetric/Hermitian generally requires pivoting
}
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
magma_csetmatrix( N, nrhs, h_B, ldb, d_B, lddb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_chesv_nopiv_gpu( opts.uplo, N, nrhs, d_A, ldda, d_B, lddb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgesv_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
//=====================================================================
// Residual
//=====================================================================
magma_cgetmatrix( N, nrhs, d_B, lddb, h_X, ldb );
Anorm = lapackf77_clange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_clange("I", &N, &nrhs, h_X, &ldb, work);
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &nrhs, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb);
Rnorm = lapackf77_clange("I", &N, &nrhs, h_B, &ldb, work);
error = Rnorm/(N*Anorm*Xnorm);
status += ! (error < tol);
/* ====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lwork = -1;
lapackf77_chesv( lapack_uplo_const(opts.uplo), &N,&nrhs,
h_A, &lda, ipiv, h_B, &ldb, &temp, &lwork, &info );
lwork = (magma_int_t) MAGMA_C_REAL( temp );
TESTING_MALLOC_PIN( hwork, magmaFloatComplex, lwork );
cpu_time = magma_wtime();
lapackf77_chesv( lapack_uplo_const(opts.uplo), &N, &nrhs,
h_A, &lda, ipiv, h_B, &ldb, hwork, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_chesv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) nrhs, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
TESTING_FREE_CPU( hwork );
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) nrhs, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A, *h_R, *work;
magmaFloatComplex *d_A, *dwork;
magmaFloatComplex c_neg_one = MAGMA_C_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};
magmaFloatComplex 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_cgetri_nb( N );
gflops = FLOPS_CGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0)
printf("lapackf77_cgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = int( MAGMA_C_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, ldwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
error = lapackf77_clange( "f", &N, &N, h_A, &lda, rwork ); // norm(A)
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
magma_cgetrf_gpu( N, N, d_A, ldda, ipiv, &info );
magma_cgetmatrix( N, N, d_A, ldda, h_A, lda );
if ( info != 0 )
printf("magma_cgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_C_MAKE( 0.0, 0.0 );
//magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgetri_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_cgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_clange( "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;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf_mgpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex *d_lA[ MagmaMaxGPUs ];
magma_int_t M, N, n2, lda, ldda, n_local, ngpu;
magma_int_t info, min_mn, nb, lhwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= (opts.check == 2); // check (-c2) implies lapack (-l)
magma_int_t status = 0;
float tol, eps = lapackf77_slamch("E");
tol = opts.tolerance * eps;
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 /(M*||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;
ldda = ((M+31)/32)*32;
nb = magma_get_cgeqrf_nb( M );
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
// ngpu must be at least the number of blocks
ngpu = min( opts.ngpu, int((N+nb-1)/nb) );
if ( ngpu < opts.ngpu ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
}
// query for workspace size
lhwork = -1;
lapackf77_cgeqrf( &M, &N, NULL, &M, NULL, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_C_REAL( tmp[0] );
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lhwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
// Allocate device memory
for( int dev = 0; dev < ngpu; dev++ ) {
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaFloatComplex, ldda*n_local );
}
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
magmaFloatComplex *tau2;
TESTING_MALLOC_CPU( tau2, magmaFloatComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_cgeqrf( &M, &N, h_A, &M, tau2, h_work, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau2 );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix_1D_col_bcyclic( M, N, h_R, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_cgeqrf2_mgpu( ngpu, M, N, d_lA, ldda, tau, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgeqrf2 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_R, lda, ngpu, nb );
magma_queue_sync( NULL );
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
magmaFloatComplex *h_W1, *h_W2, *h_W3;
float *h_RW, results[2];
TESTING_MALLOC_CPU( h_W1, magmaFloatComplex, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, magmaFloatComplex, n2 ); // R
TESTING_MALLOC_CPU( h_W3, magmaFloatComplex, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, float, M ); // RWORK
lapackf77_clarnv( &ione, ISEED2, &n2, h_A );
lapackf77_cqrt02( &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_clange("f", &M, &N, h_A, &lda, work );
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_clange("f", &M, &N, h_R, &lda, work ) / (min_mn*error);
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 );
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 );
for( int dev=0; dev < ngpu; dev++ ){
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgelqf_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_A;
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};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||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;
nb = magma_get_zgeqrf_nb(M);
gflops = FLOPS_ZGELQF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgelqf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, M*nb );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, lda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, N, h_R, lda, d_A, lda );
gpu_time = magma_wtime();
magma_zgelqf_gpu( M, N, d_A, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgelqf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgelqf(&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("lapack_zgelqf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( M, N, d_A, lda, h_R, lda );
error = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("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, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhetrd_he2hb
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, gpu_perf, gflops;
magmaDoubleComplex *h_A, *h_R, *h_work, *dT1;
magmaDoubleComplex *tau;
double *D, *E;
/* Matrix size */
magma_int_t N, n2, lda, lwork, lwork0; //ldt
magma_int_t info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
#if defined(CHECKEIG)
#if defined(PRECISION_z) || defined(PRECISION_d)
magma_int_t WANTZ=0;
magma_int_t THREADS=1;
#endif
#endif
magma_int_t NE = 0;
magma_int_t NB = 0;
magma_int_t ngpu = 1;
magma_opts opts;
opts.parse_opts( argc, argv );
NB = opts.nb;
if (NB < 1)
NB = 64; //64; //magma_get_zhetrd_he2hb_nb(N);
// what is NE ?
if (NE < 1)
NE = 64; //N; //magma_get_zhetrd_he2hb_nb(N); // N not yet initialized
printf("%% N GPU GFlop/s \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;
//ldt = N;
n2 = lda*N;
gflops = FLOPS_ZHETRD( N ) / 1e9;
/* We suppose the magma NB is bigger than lapack NB */
lwork0 = N*NB;
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, N-1 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork0 );
TESTING_MALLOC_PIN( D, double, N );
TESTING_MALLOC_PIN( E, double, N );
//TESTING_MALLOC_DEV( dT1, magmaDoubleComplex, (2*min(N,N) + roundup( N, 32 ))*NB );
TESTING_MALLOC_DEV( dT1, magmaDoubleComplex, (N*NB) );
// if (WANTZ) gflops = 2.0*gflops;
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hermitian( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_device_t cdev;
magma_getdevice( &cdev );
gpu_time = magma_wtime();
/*
magma_zhetrd_he2hb( opts.uplo, N, NB, h_R, lda, tau, h_work, lwork0, dT1, THREADS, &info);
tband = magma_wtime - gpu_time();
printf(" Finish BAND N %d NB %d ngpu %d timing= %f\n", N, NB, ngpu, tband);
magma_zhetrd_bhe2trc_v5(THREADS, WANTZ, opts.uplo, NE, N, NB, h_R, lda, D, E, dT1, ldt);
*/
/*
magma_zhetrd_he2hb( opts.uplo, N, NB, h_R, lda, tau, h_work, lwork, dT1, THREADS, &info);
tband = magma_wtime - gpu_time();
printf(" Finish BAND N %d NB %d ngpu %d timing= %f\n", N, NB, ngpu, tband);
magma_zhetrd_bhe2trc(THREADS, WANTZ, opts.uplo, NE, N, NB, h_R, lda, D, E, dT1, ldt);
*/
magma_range_t range = MagmaRangeAll;
magma_int_t m1 = 0;
double vl = 0;
double 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) );
}
magmaDoubleComplex *hh_work;
magma_int_t *iwork;
magma_int_t /*nb,*/ /*lwork,*/ liwork;
magma_int_t threads = magma_get_parallel_numthreads();
#ifdef COMPLEX
double *rwork;
magma_int_t lrwork;
#endif
magma_zheevdx_getworksize(N, threads, (opts.jobz == MagmaVec),
&lwork,
#ifdef COMPLEX
&lrwork,
#endif
&liwork);
TESTING_MALLOC_PIN( hh_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
#ifdef COMPLEX
TESTING_MALLOC_PIN( rwork, double, lrwork );
#endif
if (ngpu == 1) {
printf("calling zheevdx_2stage 1 GPU\n");
magma_zheevdx_2stage( opts.jobz, range, opts.uplo, N,
h_R, lda,
vl, vu, il, iu,
&m1, D,
hh_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info);
} else {
printf("calling zheevdx_2stage_m %d GPU\n", (int) ngpu);
magma_zheevdx_2stage_m(ngpu, opts.jobz, range, opts.uplo, N,
h_R, lda,
vl, vu, il, iu,
&m1, D,
hh_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info);
}
magma_setdevice( cdev );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
/* =====================================================================
Check the factorization
=================================================================== */
/*
if ( opts.check ) {
FILE *fp;
printf("Writing input matrix in matlab_i_mat.txt ...\n");
fp = fopen ("matlab_i_mat.txt", "w");
if ( fp == NULL ) {
printf("Couldn't open output file\n");
return -1;
}
for (j=0; j < N; j++) {
for (k=0; k < N; k++) {
#ifdef COMPLEX
fprintf(fp, "%5d %5d %11.8f %11.8f\n", k+1, j+1,
h_A[k+j*lda].x, h_A[k+j*lda].y);
#else
fprintf(fp, "%5d %5d %11.8f\n", k+1, j+1, h_A[k+j*lda]);
#endif
}
}
fclose( fp );
printf("Writing output matrix in matlab_o_mat.txt ...\n");
fp = fopen ("matlab_o_mat.txt", "w");
if ( fp == NULL ) {
printf("Couldn't open output file\n");
return -1;
}
for (j=0; j < N; j++) {
for (k=0; k < N; k++) {
#ifdef COMPLEX
fprintf(fp, "%5d %5d %11.8f %11.8f\n", k+1, j+1,
h_R[k+j*lda].x, h_R[k+j*lda].y);
#else
fprintf(fp, "%5d %5d %11.8f\n", k+1, j+1, h_R[k+j*lda]);
#endif
}
}
fclose( fp );
}
*/
/* =====================================================================
Print performance and error.
=================================================================== */
#if defined(CHECKEIG)
#if defined(PRECISION_z) || defined(PRECISION_d)
if ( opts.check ) {
printf(" Total N %5d gflops %6.2f timing %6.2f seconds\n", (int) N, gpu_perf, gpu_time );
const char* JOBZ;
if (WANTZ == 0)
JOBZ = MagmaNoVecStr;
else
JOBZ = MagmaVecStr;
double nrmI=0.0, nrm1=0.0, nrm2=0.0;
int lwork2 = 256*N;
magmaDoubleComplex *work2, *AINIT;
double *rwork2, *D2;
// TODO free this memory !
magma_zmalloc_cpu( &work2, lwork2 );
magma_dmalloc_cpu( &rwork2, N );
magma_dmalloc_cpu( &D2, N );
magma_zmalloc_cpu( &AINIT, N*lda );
memcpy(AINIT, h_A, N*lda*sizeof(magmaDoubleComplex));
/* compute the eigenvalues using lapack routine to be able to compare to it and used as ref */
cpu_time = magma_wtime();
i= min(12, THREADS);
magma_set_lapack_numthreads( i );
lapackf77_zheev( "N", "L", &N, h_A, &lda, D2, work2, &lwork2,
#ifdef COMPLEX
rwork2,
#endif
&info );
///* call eigensolver for our resulting tridiag [D E] and for Q */
//dstedc_withZ('V', N, D, E, h_R, lda);
////dsterf_( &N, D, E, &info);
////
cpu_time = magma_wtime() - cpu_time;
printf(" Finish CHECK - EIGEN timing= %f threads %d\n", cpu_time, i);
/*
for (i=0; i < 10; i++)
printf(" voici lpk D[%d] %8.2e\n", i, D2[i]);
*/
//magmaDoubleComplex mydz=0.0, mydo=1.0;
//magmaDoubleComplex *Z;
// magma_zmalloc_cpu( &Z, N*lda );
// dgemm_("N", "N", &N, &N, &N, &mydo, h_R, &lda, h_A, &lda, &mydz, Z, &lda);
/* compare result */
cmp_vals(N, D2, D, &nrmI, &nrm1, &nrm2);
magmaDoubleComplex *WORKAJETER;
double *RWORKAJETER, *RESU;
// TODO free this memory !
magma_zmalloc_cpu( &WORKAJETER, (2* N * N + N) );
magma_dmalloc_cpu( &RWORKAJETER, N );
magma_dmalloc_cpu( &RESU, 10 );
int MATYPE;
memset(RESU, 0, 10*sizeof(double));
MATYPE=3;
double NOTHING=0.0;
cpu_time = magma_wtime();
// check results
zcheck_eig_( JOBZ, &MATYPE, &N, &NB, AINIT, &lda, &NOTHING, &NOTHING, D2, D, h_R, &lda, WORKAJETER, RWORKAJETER, RESU );
cpu_time = magma_wtime() - cpu_time;
printf(" Finish CHECK - results timing= %f\n", cpu_time);
magma_set_lapack_numthreads( 1 );
printf("\n");
printf(" ================================================================================================================\n");
printf(" ==> INFO voici threads=%d N=%d NB=%d WANTZ=%d\n", (int) THREADS, (int) N, (int) NB, (int) WANTZ);
printf(" ================================================================================================================\n");
printf(" DSBTRD : %15s \n", "STATblgv9withQ ");
printf(" ================================================================================================================\n");
if (WANTZ > 0)
printf(" | A - U S U' | / ( |A| n ulp ) : %15.3E \n", RESU[0]);
if (WANTZ > 0)
printf(" | I - U U' | / ( n ulp ) : %15.3E \n", RESU[1]);
printf(" | D1 - EVEIGS | / (|D| ulp) : %15.3E \n", RESU[2]);
printf(" max | D1 - EVEIGS | : %15.3E \n", RESU[6]);
printf(" ================================================================================================================\n\n\n");
printf(" ****************************************************************************************************************\n");
printf(" * Hello here are the norm Infinite (max)=%8.2e norm one (sum)=%8.2e norm2(sqrt)=%8.2e *\n", nrmI, nrm1, nrm2);
printf(" ****************************************************************************************************************\n\n");
}
#endif
#endif
printf(" Total N %5d gflops %6.2f timing %6.2f seconds\n", (int) N, gpu_perf, gpu_time );
printf("%%===========================================================================\n\n\n");
/* Memory clean up */
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( tau );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( D );
TESTING_FREE_PIN( E );
TESTING_FREE_DEV( dT1 );
/* TODO - not all memory has been freed inside loop */
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return EXIT_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrtri
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double Anorm, error, work[1];
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% N CPU Gflop/s (sec) GPU Gflop/s (sec) ||R||_F / ||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;
gflops = FLOPS_ZTRTRI( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
if ( opts.warmup ) {
magma_zpotrf( opts.uplo, N, h_R, lda, &info );
magma_ztrtri( opts.uplo, opts.diag, N, h_R, lda, &info );
lapackf77_zlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
}
/* factorize matrix */
magma_zpotrf( opts.uplo, N, h_R, lda, &info );
// check for exact singularity
//h_R[ 10 + 10*lda ] = MAGMA_Z_ZERO;
gpu_time = magma_wtime();
magma_ztrtri( opts.uplo, opts.diag, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_ztrtri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_ztrtri( lapack_uplo_const(opts.uplo), lapack_diag_const(opts.diag), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_ztrtri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
Anorm = lapackf77_zlantr("f", lapack_uplo_const(opts.uplo), MagmaNonUnitStr, &N, &N, h_A, &lda, work);
error = lapackf77_zlantr("f", lapack_uplo_const(opts.uplo), MagmaNonUnitStr, &N, &N, h_R, &lda, work) / Anorm;
bool okay = (error < tol);
status += ! okay;
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, (okay ? "ok" : "failed") );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhetrd_he2hb
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
magma_timestr_t start, end;
double eps, flops, gpu_perf, gpu_time;
magmaDoubleComplex *h_A, *h_R, *h_work;
magmaDoubleComplex *tau;
double *D, *E;
/* Matrix size */
magma_int_t N = 0, n2, lda, lwork,ldt;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info, checkres, once = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
char *uplo = (char *)MagmaLowerStr;
magma_int_t ngpu = magma_num_gpus();
magma_int_t nstream = max(3,ngpu+1);
magma_int_t WANTZ=0;
magma_int_t THREADS=1;
magma_int_t NE = 0;
magma_int_t NB = 0;
magma_int_t distblk =0;
magma_int_t ver =0;
checkres = 0; //getenv("MAGMA_TESTINGS_CHECK") != NULL;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0) {
N = atoi(argv[++i]);
once = 1;
}
else if (strcmp("-NB", argv[i])==0) {
NB = atoi(argv[++i]);
}
else if (strcmp("-D", argv[i])==0) {
distblk = atoi(argv[++i]);
}
else if (strcmp("-threads", argv[i])==0) {
THREADS = atoi(argv[++i]);
}
else if (strcmp("-wantz", argv[i])==0) {
WANTZ = atoi(argv[++i]);
}
else if (strcmp("-NE", argv[i])==0) {
NE = atoi(argv[++i]);
}
else if ( strcmp("-c", argv[i]) == 0 ) {
checkres = 1;
}
else if ( strcmp("-v", argv[i]) == 0 && i+1 < argc ) {
ver = atoi( argv[++i] );
}
else if ( strcmp("-nstream", argv[i]) == 0 && i+1 < argc ) {
nstream = atoi( argv[++i] );
magma_assert( nstream > 0 && nstream <= 20,
"error: -nstream %s is invalid; must be > 0 and <= 20.\n", argv[i] );
}
else if ( strcmp("-ngpu", argv[i]) == 0 && i+1 < argc ) {
ngpu = atoi( argv[++i] );
magma_assert( ngpu > 0 || ngpu > MagmaMaxGPUs, "error: -ngpu %s is invalid; must be > 0.\n", argv[i] );
}
else if (strcmp("-U", argv[i])==0)
uplo = (char *)MagmaUpperStr;
else if (strcmp("-L", argv[i])==0)
uplo = (char *)MagmaLowerStr;
}
if ( N > 0 )
printf(" testing_zhetrd_he2hb -L|U -N %d -NB %d -wantz %d -threads %d check %d \n\n", N, NB, WANTZ, THREADS, checkres);
else
{
printf("\nUsage: \n");
printf(" testing_zhetrd_he2hb -L|U -N %d -NB -wantz -threads \n\n", 1024);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_zhetrd_he2hb -L|U -N %d\n\n", 1024);
N = size[9];
}
eps = lapackf77_dlamch( "E" );
lda = N;
ldt = N;
n2 = lda * N;
if(NB<1)
NB = 64; //64; //magma_get_zhetrd_he2hb_nb(N);
if(NE<1)
NE = N; //64; //magma_get_zhetrd_he2hb_nb(N);
/* We suppose the magma NB is bigger than lapack NB */
lwork = N*NB;
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, N-1 );
TESTING_MALLOC_PIN( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_PIN( D, double, N );
TESTING_MALLOC_PIN( E, double, N );
nstream = max(3,ngpu+2);
magma_queue_t streams[MagmaMaxGPUs][20];
magmaDoubleComplex *da[MagmaMaxGPUs],*dT1[MagmaMaxGPUs];
magma_int_t ldda = ((N+31)/32)*32;
if((distblk==0)||(distblk<NB))
distblk = max(256,NB);
printf("voici ngpu %d distblk %d NB %d nstream %d\n ",ngpu,distblk,NB,nstream);
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t mlocal = ((N / distblk) / ngpu + 1) * distblk;
magma_setdevice( dev );
TESTING_MALLOC_DEV( da[dev], magmaDoubleComplex, ldda*mlocal );
TESTING_MALLOC_DEV( dT1[dev], magmaDoubleComplex, N*NB );
for( int i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[dev][i] );
}
}
magma_setdevice( 0 );
for(i=0; i<10; i++){
if ( !once ) {
N = size[i];
}
lda = N;
n2 = N*lda;
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hermitian( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* Copy the matrix to the GPU */
magma_zsetmatrix_1D_col_bcyclic( N, N, h_R, lda, da, ldda, ngpu, distblk);
//magmaDoubleComplex *dabis;
// TESTING_MALLOC_DEV( dabis, magmaDoubleComplex, ldda*N );
// magma_zsetmatrix(N,N,h_R,lda,dabis,ldda);
for (int count=0; count<1;++count){
magma_setdevice(0);
start = get_current_time();
if(ver==30){
magma_zhetrd_he2hb_mgpu_spec(uplo[0], N, NB, h_R, lda, tau, h_work, lwork, da, ldda, dT1, NB, ngpu, distblk, streams, nstream, THREADS, &info);
}else{
nstream =3;
magma_zhetrd_he2hb_mgpu(uplo[0], N, NB, h_R, lda, tau, h_work, lwork, da, ldda, dT1, NB, ngpu, distblk, streams, nstream, THREADS, &info);
}
// magma_zhetrd_he2hb(uplo[0], N, NB, h_R, lda, tau, h_work, lwork, dT1[0], &info);
end = get_current_time();
printf(" Finish BAND N %d NB %d dist %d ngpu %d version %d timing= %f\n", N, NB, distblk, ngpu, ver, GetTimerValue(start,end) / 1000.);
}
magma_setdevice(0);
//goto fin;
//return 0;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice(dev);
cudaDeviceSynchronize();
}
magma_setdevice(0);
magmablasSetKernelStream( NULL );
magma_zhetrd_bhe2trc_v5(THREADS, WANTZ, uplo[0], NE, N, NB, h_R, lda, D, E, dT1[0], ldt);
// magma_zhetrd_bhe2trc(THREADS, WANTZ, uplo[0], NE, N, NB, h_R, lda, D, E, dT1[0], ldt);
end = get_current_time();
if (info != 0)
printf("magma_zhetrd_he2hb returned error %d: %s.\n",
(int) info, magma_strerror( info ));
gpu_perf = flops / GetTimerValue(start,end);
gpu_time = GetTimerValue(start,end) / 1000.;
/* =====================================================================
Print performance and error.
=================================================================== */
#if defined(CHECKEIG)
#if defined(PRECISION_z) || defined(PRECISION_d)
if ( checkres ) {
printf(" Total N %5d flops %6.2f timing %6.2f seconds\n", (int) N, gpu_perf, gpu_time );
char JOBZ;
if(WANTZ==0)
JOBZ='N';
else
JOBZ = 'V';
double nrmI=0.0, nrm1=0.0, nrm2=0.0;
int lwork2 = 256*N;
magmaDoubleComplex *work2 = (magmaDoubleComplex *) malloc (lwork2*sizeof(magmaDoubleComplex));
double *rwork2 = (double *) malloc (N*sizeof(double));
double *D2 = (double *) malloc (N*sizeof(double));
magmaDoubleComplex *AINIT = (magmaDoubleComplex *) malloc (N*lda*sizeof(magmaDoubleComplex));
memcpy(AINIT, h_A, N*lda*sizeof(magmaDoubleComplex));
/* compute the eigenvalues using lapack routine to be able to compare to it and used as ref */
start = get_current_time();
i= min(12,THREADS);
#if defined(USEMKL)
mkl_set_num_threads( i );
#endif
#if defined(USEACML)
omp_set_num_threads(i);
#endif
#if defined(PRECISION_z) || defined (PRECISION_c)
lapackf77_zheev( "N", "L", &N, h_A, &lda, D2, work2, &lwork2, rwork2, &info );
#else
lapackf77_dsyev( "N", "L", &N, h_A, &lda, D2, work2, &lwork2, &info );
#endif
///* call eigensolver for our resulting tridiag [D E] and for Q */
//dstedc_withZ('V', N, D, E, h_R, lda);
////dsterf_( &N, D, E, &info);
////
end = get_current_time();
printf(" Finish CHECK - EIGEN timing= %f threads %d\n", GetTimerValue(start,end) / 1000., i);
/* compare result */
cmp_vals(N, D2, D, &nrmI, &nrm1, &nrm2);
magmaDoubleComplex *WORKAJETER;
double *RWORKAJETER, *RESU;
WORKAJETER = (magmaDoubleComplex *) malloc( (2* N * N + N) * sizeof(magmaDoubleComplex) );
RWORKAJETER = (double *) malloc( N * sizeof(double) );
RESU = (double *) malloc(10*sizeof(double));
int MATYPE;
memset(RESU,0,10*sizeof(double));
MATYPE=3;
double NOTHING=0.0;
start = get_current_time();
// check results
zcheck_eig_(&JOBZ, &MATYPE, &N, &NB, AINIT, &lda, &NOTHING, &NOTHING, D2 , D, h_R, &lda, WORKAJETER, RWORKAJETER, RESU );
end = get_current_time();
printf(" Finish CHECK - results timing= %f\n", GetTimerValue(start,end) / 1000.);
#if defined(USEMKL)
mkl_set_num_threads( 1 );
#endif
#if defined(USEACML)
omp_set_num_threads(1);
#endif
printf("\n");
printf(" ================================================================================================================\n");
printf(" ==> INFO voici threads=%d N=%d NB=%d WANTZ=%d\n", (int) THREADS, (int) N, (int) NB, (int) WANTZ);
printf(" ================================================================================================================\n");
printf(" DSBTRD : %15s \n", "STATblgv9withQ ");
printf(" ================================================================================================================\n");
if(WANTZ>0)
printf(" | A - U S U' | / ( |A| n ulp ) : %15.3E \n",RESU[0]);
if(WANTZ>0)
printf(" | I - U U' | / ( n ulp ) : %15.3E \n", RESU[1]);
printf(" | D1 - EVEIGS | / (|D| ulp) : %15.3E \n", RESU[2]);
printf(" max | D1 - EVEIGS | : %15.3E \n", RESU[6]);
printf(" ================================================================================================================\n\n\n");
printf(" ****************************************************************************************************************\n");
printf(" * Hello here are the norm Infinite (max)=%8.2e norm one (sum)=%8.2e norm2(sqrt)=%8.2e *\n", nrmI, nrm1, nrm2);
printf(" ****************************************************************************************************************\n\n");
}
#endif
#endif
printf(" Total N %5d flops %6.2f timing %6.2f seconds\n", (int) N, 0.0, gpu_time );
printf("============================================================================\n\n\n");
if ( once )
break;
}
//fin:
/* Memory clean up */
TESTING_FREE_CPU( tau );
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( D );
TESTING_FREE_PIN( E );
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
TESTING_FREE_DEV( da[dev] );
TESTING_FREE_DEV( dT1[dev] );
for( int i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[dev][i] );
}
}
magma_setdevice( 0 );
/* Shutdown */
TESTING_FINALIZE_MGPU();
return EXIT_SUCCESS;
}
