/* //////////////////////////////////////////////////////////////////////////// -- Testing cungqr */ 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 *hA, *hR, *hT, *tau, *h_work; magmaFloatComplex *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 ); opts.lapack |= opts.check; // check (-c) implies lapack (-l) float tol = opts.tolerance * lapackf77_slamch("E"); 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( hA, magmaFloatComplex, lda*n ); TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork ); TESTING_MALLOC_CPU( hR, magmaFloatComplex, lda*n ); TESTING_MALLOC_CPU( hT, magmaFloatComplex, min_mn*nb ); 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( MagmaUpperLowerStr, &m, &n, hA, &lda, hR, &lda ); /* ==================================================================== Performs operation using MAGMA =================================================================== */ // first, get QR factors 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, hR, lda ); magma_cgetmatrix( nb, min_mn, dT, nb, hT, nb ); gpu_time = magma_wtime(); magma_cungqr_m( m, n, k, hR, lda, tau, hT, nb, &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 ) { error = lapackf77_clange("f", &m, &n, hA, &lda, work ); lapackf77_cgeqrf( &m, &n, hA, &lda, tau, h_work, &lwork, &info ); if ( info != 0 ) printf("lapackf77_cgeqrf returned error %d: %s.\n", (int) info, magma_strerror( info )); 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) / error; 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, (error < tol ? "ok" : "failed") ); status += ! (error < tol); } 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( hT ); 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 ctrmm */ int main( int argc, char** argv) { TESTING_INIT(); real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time; float cublas_error, Cnorm, work[1]; magma_int_t M, N; magma_int_t Ak; magma_int_t sizeA, sizeB; magma_int_t lda, ldb, ldda, lddb; magma_int_t ione = 1; magma_int_t ISEED[4] = {0,0,0,1}; magmaFloatComplex *h_A, *h_B, *h_Bcublas; magmaFloatComplex *d_A, *d_B; magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE; magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 ); 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("If running lapack (option --lapack), CUBLAS error is computed\n" "relative to CPU BLAS result.\n\n"); printf("side = %s, uplo = %s, transA = %s, diag = %s \n", lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) ); printf(" M N CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\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]; gflops = FLOPS_CTRMM(opts.side, M, N) / 1e9; if ( opts.side == MagmaLeft ) { lda = M; Ak = M; } else { lda = N; Ak = N; } ldb = M; ldda = ((lda+31)/32)*32; lddb = ((ldb+31)/32)*32; sizeA = lda*Ak; sizeB = ldb*N; TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*Ak ); TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*N ); TESTING_MALLOC_CPU( h_Bcublas, magmaFloatComplex, ldb*N ); TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*Ak ); TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*N ); /* Initialize the matrices */ lapackf77_clarnv( &ione, ISEED, &sizeA, h_A ); lapackf77_clarnv( &ione, ISEED, &sizeB, h_B ); /* ===================================================================== Performs operation using CUBLAS =================================================================== */ magma_csetmatrix( Ak, Ak, h_A, lda, d_A, ldda ); magma_csetmatrix( M, N, h_B, ldb, d_B, lddb ); // note cublas does trmm out-of-place (i.e., adds output matrix C), // but allows C=B to do in-place. cublas_time = magma_sync_wtime( NULL ); cublasCtrmm( handle, cublas_side_const(opts.side), cublas_uplo_const(opts.uplo), cublas_trans_const(opts.transA), cublas_diag_const(opts.diag), M, N, &alpha, d_A, ldda, d_B, lddb, d_B, lddb ); cublas_time = magma_sync_wtime( NULL ) - cublas_time; cublas_perf = gflops / cublas_time; magma_cgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb ); /* ===================================================================== Performs operation using CPU BLAS =================================================================== */ if ( opts.lapack ) { cpu_time = magma_wtime(); blasf77_ctrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag), &M, &N, &alpha, h_A, &lda, h_B, &ldb ); cpu_time = magma_wtime() - cpu_time; cpu_perf = gflops / cpu_time; } /* ===================================================================== Check the result =================================================================== */ if ( opts.lapack ) { // compute relative error for both magma & cublas, relative to lapack, // |C_magma - C_lapack| / |C_lapack| Cnorm = lapackf77_clange( "M", &M, &N, h_B, &ldb, work ); blasf77_caxpy( &sizeB, &c_neg_one, h_B, &ione, h_Bcublas, &ione ); cublas_error = lapackf77_clange( "M", &M, &N, h_Bcublas, &ldb, work ) / Cnorm; printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n", (int) M, (int) N, cublas_perf, 1000.*cublas_time, cpu_perf, 1000.*cpu_time, cublas_error, (cublas_error < tol ? "ok" : "failed")); status += ! (cublas_error < tol); } else { printf("%5d %5d %7.2f (%7.2f) --- ( --- ) --- ---\n", (int) M, (int) N, cublas_perf, 1000.*cublas_time); } TESTING_FREE_CPU( h_A ); TESTING_FREE_CPU( h_B ); TESTING_FREE_CPU( h_Bcublas ); TESTING_FREE_DEV( d_A ); TESTING_FREE_DEV( d_B ); fflush( stdout ); } if ( opts.niter > 1 ) { printf( "\n" ); } } 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; }
/* //////////////////////////////////////////////////////////////////////////// -- Testing zgetrf */ 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, *work; magmaDoubleComplex *d_A, *dwork; magmaDoubleComplex c_neg_one = MAGMA_Z_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}; magmaDoubleComplex tmp; double 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 ); double tol = opts.tolerance * lapackf77_dlamch("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_zgetri_nb( N ); gflops = FLOPS_ZGETRI( N ) / 1e9; // query for workspace size lwork = -1; lapackf77_zgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info ); if (info != 0) printf("lapackf77_zgetri returned error %d: %s.\n", (int) info, magma_strerror( info )); lwork = int( MAGMA_Z_REAL( tmp )); TESTING_MALLOC_CPU( ipiv, magma_int_t, N ); TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork ); TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 ); TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 ); TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N ); TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork ); /* Initialize the matrix */ lapackf77_zlarnv( &ione, ISEED, &n2, h_A ); error = lapackf77_zlange( "f", &N, &N, h_A, &lda, rwork ); // norm(A) /* Factor the matrix. Both MAGMA and LAPACK will use this factor. */ magma_zsetmatrix( N, N, h_A, lda, d_A, ldda ); magma_zgetrf_gpu( N, N, d_A, ldda, ipiv, &info ); magma_zgetmatrix( N, N, d_A, ldda, h_A, lda ); if ( info != 0 ) printf("magma_zgetrf_gpu returned error %d: %s.\n", (int) info, magma_strerror( info )); // check for exact singularity //h_A[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 ); //magma_zsetmatrix( N, N, h_A, lda, d_A, ldda ); /* ==================================================================== Performs operation using MAGMA =================================================================== */ gpu_time = magma_wtime(); magma_zgetri_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_zgetri_gpu returned error %d: %s.\n", (int) info, magma_strerror( info )); magma_zgetmatrix( N, N, d_A, ldda, h_R, lda ); /* ===================================================================== Performs operation using LAPACK =================================================================== */ if ( opts.lapack ) { cpu_time = magma_wtime(); lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info ); cpu_time = magma_wtime() - cpu_time; cpu_perf = gflops / cpu_time; if (info != 0) printf("lapackf77_zgetri 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 ); error = lapackf77_zlange( "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 magma_dsymm_mgpu */ int main( int argc, char** argv) { TESTING_INIT(); double c_neg_one = MAGMA_D_NEG_ONE; double alpha = MAGMA_D_MAKE( 3.456, 5.678 ); double beta = MAGMA_D_MAKE( 1.234, 2.456 ); real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.; real_Double_t gpu_perf2=0., gpu_time2=0.; double Anorm, error, work[1]; double *hA, *hB, *hC, *hR; magmaDouble_ptr dA[MagmaMaxGPUs], dB[MagmaMaxGPUs], dC[MagmaMaxGPUs], dwork[MagmaMaxGPUs]; magmaDouble_ptr dA2; magma_int_t i, j, dev, M, N, size, lda, ldb, ldc, ldda, lddb, lddc, msize, nb; magma_int_t ione = 1; magma_int_t iseed[4] = {0,0,0,1}; magma_int_t status = 0; magma_opts opts; opts.parse_opts( argc, argv ); opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code double tol = opts.tolerance * lapackf77_dlamch("E"); // default values nb = (opts.nb > 0 ? opts.nb : 64); magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2]; magma_int_t ncmplx = 0; magma_buildconnection_mgpu( gnode, &ncmplx, opts.ngpu ); printf("%% Initializing communication pattern... GPU-ncmplx %d\n", (int) ncmplx); for (i=0; i < ncmplx; ++i) { magma_int_t myngpu = gnode[i][MagmaMaxGPUs]; printf("%% cmplx %d has %d GPUs:", i, myngpu); for (j=0; j < myngpu; ++j) { printf(" %d", (int) gnode[i][j]); if (j < myngpu-1) { printf(","); } } printf("\n"); } // number of queues per GPU. Requires ngpu. magma_int_t nqueue = opts.ngpu; // number of events per GPU. Require ngpu*ngpu. magma_int_t nevents = opts.ngpu*opts.ngpu; magma_queue_t queues[MagmaMaxGPUs][20], queues0[MagmaMaxGPUs]; magma_event_t events[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs + 10]; for( dev = 0; dev < opts.ngpu; ++dev ) { magma_setdevice( dev ); for( i = 0; i < nqueue; ++i ) { magma_queue_create( dev, &queues[dev][i] ); } queues0[dev] = queues[dev][0]; for( i = 0; i < nevents; ++i ) { cudaEventCreateWithFlags( &events[dev][i], cudaEventDisableTiming ); } } printf("%% nb %d, ngpu %d, version %d\n", (int) nb, (int) opts.ngpu, (int) opts.version ); printf("%% M N nb offset CPU Gflop/s (sec) GPU Gflop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||B||\n"); printf("%%========================================================================================================\n"); for( int itest = 0; itest < opts.ntest; ++itest ) { M = opts.msize[itest]; N = opts.nsize[itest]; for( int offset = 0; offset < N; offset += min(N,nb) ) { for( int iter = 0; iter < opts.niter; ++iter ) { msize = M - offset; lda = M; // TODO depends on side ldb = M; ldc = M; ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default lddc = magma_roundup( ldc, opts.align ); // multiple of 32 by default gflops = FLOPS_DSYMM( MagmaLeft, (double)msize, (double)N ) / 1e9; magma_int_t dworksiz = lddc*N + (M*N)*opts.ngpu; TESTING_MALLOC_CPU( hA, double, lda*M ); TESTING_MALLOC_CPU( hB, double, ldb*N ); TESTING_MALLOC_CPU( hC, double, ldc*N ); TESTING_MALLOC_PIN( hR, double, ldc*N ); for( dev = 0; dev < opts.ngpu; ++dev ) { magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb; magma_setdevice( dev ); TESTING_MALLOC_DEV( dA[dev], double, ldda*mlocal ); TESTING_MALLOC_DEV( dB[dev], double, lddb*N ); TESTING_MALLOC_DEV( dC[dev], double, lddc*N ); TESTING_MALLOC_DEV( dwork[dev], double, dworksiz ); } if ( opts.check ) { magma_setdevice( 0 ); TESTING_MALLOC_DEV( dA2, double, ldda*M ); } size = lda*M; lapackf77_dlarnv( &ione, iseed, &size, hA ); magma_dmake_symmetric( M, hA, lda ); size = ldb*N; lapackf77_dlarnv( &ione, iseed, &size, hB ); size = ldc*N; lapackf77_dlarnv( &ione, iseed, &size, hC ); lapackf77_dlacpy( "Full", &M, &N, hC, &ldc, hR, &lda ); /* ==================================================================== Performs operation using MAGMA =================================================================== */ magma_dsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb, queues0 ); for( dev = 0; dev < opts.ngpu; ++dev ) { magma_setdevice( dev ); magma_dsetmatrix( M, N, hB, lda, dB[dev], ldda, opts.queue ); // since when offset != 0, the GPU that does beta*C may not be 0, // send initial hC to all GPUs. magma_dsetmatrix( M, N, hC, lda, dC[dev], ldda, opts.queue ); } trace_init( 1, opts.ngpu, nqueue, (magma_queue_t*) queues ); gpu_time = magma_sync_wtime(0); magmablas_dsymm_mgpu( MagmaLeft, MagmaLower, msize, N, alpha, dA, ldda, offset, dB, ldda, beta, dC, ldda, dwork, dworksiz, opts.ngpu, nb, queues, nqueue, events, nevents, gnode, ncmplx); gpu_time = magma_sync_wtime(0) - gpu_time; gpu_perf = gflops / gpu_time; #ifdef TRACING char buf[80]; snprintf( buf, sizeof(buf), "dsymm-m%d-n%d-nb%d-ngpu%d-run%d.svg", (int) M, (int) N, (int) nb, (int) opts.ngpu, (int) iter ); trace_finalize( buf, "trace.css" ); #endif /* ==================================================================== Performs operation using CUBLAS =================================================================== */ if ( opts.check && iter == 0 ) { magma_setdevice( 0 ); magma_dsetmatrix( M, M, hA, lda, dA2, ldda, opts.queue ); magma_dsetmatrix( M, N, hB, lda, dB[0], ldda, opts.queue ); magma_dsetmatrix( M, N, hC, lda, dwork[0], ldda, opts.queue ); gpu_time2 = magma_sync_wtime(0); magma_dsymm( MagmaLeft, MagmaLower, msize, N, alpha, dA2 + offset + offset*ldda, ldda, dB[0], ldda, beta, dwork[0], ldda, opts.queue ); gpu_time2 = magma_sync_wtime(0) - gpu_time2; gpu_perf2 = gflops / gpu_time2; } /* ===================================================================== Performs operation using LAPACK =================================================================== */ if ( opts.check ) { // store ||A||*||B|| Anorm = lapackf77_dlange("fro", &msize, &msize, hA + offset + offset*lda, &lda, work ); Anorm *= lapackf77_dlange("fro", &msize, &N, hB, &lda, work ); //printf( "A =" ); magma_dprint( M, M, hA, lda ); //printf( "B =" ); magma_dprint( M, N, hB, lda ); //printf( "C =" ); magma_dprint( M, N, hC, lda ); cpu_time = magma_wtime(); blasf77_dsymm( "Left", "Lower", &msize, &N, &alpha, hA + offset + offset*lda, &lda, hB, &lda, &beta, hC, &lda ); cpu_time = magma_wtime() - cpu_time; cpu_perf = gflops / cpu_time; for (dev=0; dev < opts.ngpu; ++dev) { magma_setdevice( dev ); magma_dgetmatrix( M, N, dC[dev], ldda, hR, lda, opts.queue ); // compute relative error ||R||/||A||*||B||, where R := C_magma - C_lapack = R - C size = ldc*N; blasf77_daxpy( &size, &c_neg_one, hC, &ione, hR, &ione ); error = lapackf77_dlange("fro", &msize, &N, hR, &lda, work) / Anorm; //printf( "R =" ); magma_dprint( M, N, hR, lda ); bool okay = (error < tol); status += ! okay; if (dev == 0) { printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n", (int) M, (int) N, (int) nb, (int) offset, cpu_perf, cpu_time, gpu_perf, gpu_time, gpu_perf2, gpu_time2, error, (okay ? "ok" : "failed") ); } else { printf( " dev %d %74s %8.2e %s\n", dev, "", error, (okay ? "ok" : "failed") ); } } } else { printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n", (int) M, (int) N, (int) nb, (int) offset, gpu_perf, gpu_time ); } TESTING_FREE_CPU( hA ); TESTING_FREE_CPU( hB ); TESTING_FREE_CPU( hC ); TESTING_FREE_PIN( hR ); for( dev = 0; dev < opts.ngpu; ++dev ) { magma_setdevice( dev ); TESTING_FREE_DEV( dA[dev] ); TESTING_FREE_DEV( dB[dev] ); TESTING_FREE_DEV( dC[dev] ); TESTING_FREE_DEV( dwork[dev] ); } if ( opts.check ) { magma_setdevice( 0 ); TESTING_FREE_DEV( dA2 ); } fflush( stdout ); } if ( opts.niter > 1 ) { printf( "\n" ); } } // offset printf( "\n" ); } for( dev = 0; dev < opts.ngpu; ++dev ) { magma_setdevice( dev ); for( i = 0; i < nqueue; ++i ) { magma_queue_destroy( queues[dev][i] ); } for( i = 0; i < nevents; ++i ) { magma_event_destroy( events[dev][i] ); } } opts.cleanup(); TESTING_FINALIZE(); return status; }
/* //////////////////////////////////////////////////////////////////////////// -- Testing dlaset Code is very similar to testing_dlacpy.cpp */ int main( int argc, char** argv) { TESTING_INIT(); real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time; double error, work[1]; double c_neg_one = MAGMA_D_NEG_ONE; double *h_A, *h_R; magmaDouble_ptr d_A; double offdiag, diag; magma_int_t M, N, size, lda, ldda; magma_int_t ione = 1; magma_int_t status = 0; magma_opts opts; parse_opts( argc, argv, &opts ); magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull }; printf("uplo M N offdiag diag CPU GByte/s (ms) GPU GByte/s (ms) check\n"); printf("================================================================================\n"); for( int iuplo = 0; iuplo < 3; ++iuplo ) { for( int itest = 0; itest < opts.ntest; ++itest ) { for( int iter = 0; iter < opts.niter; ++iter ) { for( int ival = 0; ival < 4; ++ival ) { // test combinations of zero & non-zero: // ival offdiag diag // 0 0 0 // 1 0 3.14 // 2 1.23 0 // 3 1.23 3.14 offdiag = MAGMA_D_MAKE( 1.2345, 6.7890 ) * (ival / 2); diag = MAGMA_D_MAKE( 3.1415, 2.7183 ) * (ival % 2); M = opts.msize[itest]; N = opts.nsize[itest]; //M += 2; // space for insets //N += 2; lda = M; ldda = roundup( M, opts.roundup ); size = lda*N; if ( uplo[iuplo] == MagmaLower || uplo[iuplo] == MagmaUpper ) { // save triangle (with diagonal) // TODO wrong for trapezoid gbytes = sizeof(double) * 0.5*N*(N+1) / 1e9; } else { // save entire matrix gbytes = sizeof(double) * 1.*M*N / 1e9; } TESTING_MALLOC_CPU( h_A, double, size ); TESTING_MALLOC_CPU( h_R, double, size ); TESTING_MALLOC_DEV( d_A, double, ldda*N ); /* Initialize the matrix */ for( int j = 0; j < N; ++j ) { for( int i = 0; i < M; ++i ) { h_A[i + j*lda] = MAGMA_D_MAKE( i + j/10000., j ); } } /* ==================================================================== Performs operation using MAGMA =================================================================== */ magma_dsetmatrix( M, N, h_A, lda, d_A, ldda ); gpu_time = magma_sync_wtime( 0 ); //magmablas_dlaset( uplo[iuplo], M-2, N-2, offdiag, diag, d_A+1+ldda, ldda ); // inset by 1 row & col magmablas_dlaset( uplo[iuplo], M, N, offdiag, diag, d_A, ldda ); gpu_time = magma_sync_wtime( 0 ) - gpu_time; gpu_perf = gbytes / gpu_time; /* ===================================================================== Performs operation using LAPACK =================================================================== */ cpu_time = magma_wtime(); //magma_int_t M2 = M-2; // inset by 1 row & col //magma_int_t N2 = N-2; //lapackf77_dlaset( lapack_uplo_const( uplo[iuplo] ), &M2, &N2, &offdiag, &diag, h_A+1+lda, &lda ); lapackf77_dlaset( lapack_uplo_const( uplo[iuplo] ), &M, &N, &offdiag, &diag, h_A, &lda ); cpu_time = magma_wtime() - cpu_time; cpu_perf = gbytes / cpu_time; if ( opts.verbose ) { printf( "A= " ); magma_dprint( M, N, h_A, lda ); printf( "dA=" ); magma_dprint_gpu( M, N, d_A, ldda ); } /* ===================================================================== Check the result =================================================================== */ magma_dgetmatrix( M, N, d_A, ldda, h_R, lda ); blasf77_daxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione); error = lapackf77_dlange("f", &M, &N, h_R, &lda, work); bool okay = (error == 0); status += ! okay; printf("%5s %5d %5d %7.2f %4.2f %7.2f (%7.2f) %7.2f (%7.2f) %s\n", lapack_uplo_const( uplo[iuplo] ), (int) M, (int) N, real(offdiag), real(diag), cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., (okay ? "ok" : "failed") ); TESTING_FREE_CPU( h_A ); TESTING_FREE_CPU( h_R ); TESTING_FREE_DEV( d_A ); fflush( stdout ); } } if ( opts.niter > 1 ) { printf( "\n" ); } } printf( "\n" ); } TESTING_FINALIZE(); return status; }
/* //////////////////////////////////////////////////////////////////////////// -- Testing ztranspose Code is very similar to testing_zsymmetrize.cpp */ int main( int argc, char** argv) { TESTING_INIT(); real_Double_t gbytes, gpu_perf, gpu_time, gpu_perf2=0, gpu_time2=0, cpu_perf, cpu_time; double error, error2, work[1]; magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE; magmaDoubleComplex *h_A, *h_B, *h_R; magmaDoubleComplex_ptr d_A, d_B; magma_int_t M, N, size, lda, ldda, ldb, lddb; magma_int_t ione = 1; magma_int_t status = 0; magma_opts opts; parse_opts( argc, argv, &opts ); printf("Inplace transpose requires M==N.\n"); printf(" M N CPU GByte/s (ms) GPU GByte/s (ms) check Inplace GB/s (ms) check\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]; lda = M; ldda = ((M+31)/32)*32; ldb = N; lddb = ((N+31)/32)*32; // load entire matrix, save entire matrix gbytes = sizeof(magmaDoubleComplex) * 2.*M*N / 1e9; TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N ); // input: M x N TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, ldb*M ); // output: N x M TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, ldb*M ); // output: N x M TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N ); // input: M x N TESTING_MALLOC_DEV( d_B, magmaDoubleComplex, lddb*M ); // output: N x M /* Initialize the matrix */ for( int j = 0; j < N; ++j ) { for( int i = 0; i < M; ++i ) { h_A[i + j*lda] = MAGMA_Z_MAKE( i + j/10000., j ); } } for( int j = 0; j < M; ++j ) { for( int i = 0; i < N; ++i ) { h_B[i + j*ldb] = MAGMA_Z_MAKE( i + j/10000., j ); } } magma_zsetmatrix( N, M, h_B, ldb, d_B, lddb ); /* ===================================================================== Performs operation using naive out-of-place algorithm (LAPACK doesn't implement transpose) =================================================================== */ cpu_time = magma_wtime(); //for( int j = 1; j < N-1; ++j ) { // inset by 1 row & col // for( int i = 1; i < M-1; ++i ) { // inset by 1 row & col for( int j = 0; j < N; ++j ) { for( int i = 0; i < M; ++i ) { h_B[j + i*ldb] = h_A[i + j*lda]; } } cpu_time = magma_wtime() - cpu_time; cpu_perf = gbytes / cpu_time; /* ==================================================================== Performs operation using MAGMA, out-of-place =================================================================== */ magma_zsetmatrix( M, N, h_A, lda, d_A, ldda ); magma_zsetmatrix( N, M, h_B, ldb, d_B, lddb ); gpu_time = magma_sync_wtime( 0 ); //magmablas_ztranspose( M-2, N-2, d_A+1+ldda, ldda, d_B+1+lddb, lddb ); // inset by 1 row & col magmablas_ztranspose( M, N, d_A, ldda, d_B, lddb ); gpu_time = magma_sync_wtime( 0 ) - gpu_time; gpu_perf = gbytes / gpu_time; /* ==================================================================== Performs operation using MAGMA, in-place =================================================================== */ if ( M == N ) { magma_zsetmatrix( M, N, h_A, lda, d_A, ldda ); gpu_time2 = magma_sync_wtime( 0 ); //magmablas_ztranspose_inplace( N-2, d_A+1+ldda, ldda ); // inset by 1 row & col magmablas_ztranspose_inplace( N, d_A, ldda ); gpu_time2 = magma_sync_wtime( 0 ) - gpu_time2; gpu_perf2 = gbytes / gpu_time2; } /* ===================================================================== Check the result =================================================================== */ // check out-of-place transpose (d_B) size = ldb*M; magma_zgetmatrix( N, M, d_B, lddb, h_R, ldb ); blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione ); error = lapackf77_zlange("f", &N, &M, h_R, &ldb, work ); if ( M == N ) { // also check in-place tranpose (d_A) magma_zgetmatrix( N, M, d_A, ldda, h_R, ldb ); blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione ); error2 = lapackf77_zlange("f", &N, &M, h_R, &ldb, work ); printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %6s %7.2f (%7.2f) %s\n", (int) M, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., (error == 0. ? "ok" : "failed"), gpu_perf2, gpu_time2, (error2 == 0. ? "ok" : "failed") ); status += ! (error == 0. && error2 == 0.); } else { printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %6s --- ( --- )\n", (int) M, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., (error == 0. ? "ok" : "failed") ); status += ! (error == 0.); } TESTING_FREE_CPU( h_A ); TESTING_FREE_CPU( h_B ); TESTING_FREE_CPU( h_R ); TESTING_FREE_DEV( d_A ); TESTING_FREE_DEV( d_B ); fflush( stdout ); } if ( opts.niter > 1 ) { printf( "\n" ); } } TESTING_FINALIZE(); return status; }
/* //////////////////////////////////////////////////////////////////////////// -- Testing slacpy */ int main( int argc, char** argv) { TESTING_INIT(); real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time; float error, work[1]; float c_neg_one = MAGMA_S_NEG_ONE; float *h_A, *h_B, *h_R; float *d_A, *d_B; magma_int_t M, N, size, lda, ldb, ldda, lddb; magma_int_t ione = 1; magma_int_t status = 0; magma_opts opts; parse_opts( argc, argv, &opts ); magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull }; printf("uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n"); printf("=================================================================\n"); for( int iuplo = 0; iuplo < 3; ++iuplo ) { for( int itest = 0; itest < opts.ntest; ++itest ) { for( int iter = 0; iter < opts.niter; ++iter ) { M = opts.msize[itest]; N = opts.nsize[itest]; lda = M; ldb = lda; ldda = ((M+31)/32)*32; lddb = ldda; size = lda*N; if ( uplo[iuplo] == MagmaLower || uplo[iuplo] == MagmaUpper ) { // load and save triangle (with diagonal) gbytes = sizeof(float) * 1.*N*(N+1) / 1e9; } else { // load entire matrix, save entire matrix gbytes = sizeof(float) * 2.*M*N / 1e9; } TESTING_MALLOC_CPU( h_A, float, size ); TESTING_MALLOC_CPU( h_B, float, size ); TESTING_MALLOC_CPU( h_R, float, size ); TESTING_MALLOC_DEV( d_A, float, ldda*N ); TESTING_MALLOC_DEV( d_B, float, ldda*N ); /* Initialize the matrix */ for( int j = 0; j < N; ++j ) { for( int i = 0; i < M; ++i ) { h_A[i + j*lda] = MAGMA_S_MAKE( i + j/10000., j ); h_B[i + j*ldb] = MAGMA_S_MAKE( i - j/10000. + 10000., j ); } } /* ==================================================================== Performs operation using MAGMA =================================================================== */ magma_ssetmatrix( M, N, h_A, lda, d_A, ldda ); magma_ssetmatrix( M, N, h_B, ldb, d_B, lddb ); gpu_time = magma_sync_wtime( 0 ); //magmablas_slacpy( uplo[iuplo], M-2, N-2, d_A+1+ldda, ldda, d_B+1+lddb, lddb ); // inset by 1 row & col magmablas_slacpy( uplo[iuplo], M, N, d_A, ldda, d_B, lddb ); gpu_time = magma_sync_wtime( 0 ) - gpu_time; gpu_perf = gbytes / gpu_time; /* ===================================================================== Performs operation using LAPACK =================================================================== */ cpu_time = magma_wtime(); //magma_int_t M2 = M-2; // inset by 1 row & col //magma_int_t N2 = N-2; //lapackf77_slacpy( lapack_uplo_const(uplo[iuplo]), &M2, &N2, h_A+1+lda, &lda, h_B+1+ldb, &ldb ); lapackf77_slacpy( lapack_uplo_const(uplo[iuplo]), &M, &N, h_A, &lda, h_B, &ldb ); cpu_time = magma_wtime() - cpu_time; cpu_perf = gbytes / cpu_time; if ( opts.verbose ) { printf( "A= " ); magma_sprint( M, N, h_A, lda ); printf( "B= " ); magma_sprint( M, N, h_B, lda ); printf( "dA=" ); magma_sprint_gpu( M, N, d_A, ldda ); printf( "dB=" ); magma_sprint_gpu( M, N, d_B, ldda ); } /* ===================================================================== Check the result =================================================================== */ magma_sgetmatrix( M, N, d_B, ldda, h_R, lda ); blasf77_saxpy(&size, &c_neg_one, h_B, &ione, h_R, &ione); error = lapackf77_slange("f", &M, &N, h_R, &lda, work); printf("%5s %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n", lapack_uplo_const(uplo[iuplo]), (int) M, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., (error == 0. ? "ok" : "failed") ); status += ! (error == 0.); TESTING_FREE_CPU( h_A ); TESTING_FREE_CPU( h_B ); TESTING_FREE_CPU( h_R ); TESTING_FREE_DEV( d_A ); TESTING_FREE_DEV( d_B ); fflush( stdout ); } if ( opts.niter > 1 ) { printf( "\n" ); } } printf( "\n" ); } TESTING_FINALIZE(); return status; }