int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
double *A, *B, *C, *C2, *LU;
double *dA, *dB, *dC1, *dC2;
double alpha = MAGMA_D_MAKE( 0.5, 0.1 );
double beta = MAGMA_D_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_err_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
printf("=========================================================================\n");
printf( "M %d, N %d, K %d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_dmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_dmalloc( &dA, size ); assert( err == 0 );
err = magma_dmalloc( &dB, size ); assert( err == 0 );
err = magma_dmalloc( &dC1, size ); assert( err == 0 );
err = magma_dmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_dlarnv( &ione, ISEED, &size, B );
lapackf77_dlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test DSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetmatrix( m, n, A, ld, dB, ld );
magma_dswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_dswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasDaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_dgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "dswap diff %.2g\n", error );
// ----- test IDAMAX
// get argmax of column of A
magma_dsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_idamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIdamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "idamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test DGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetvector( maxn, B, 1, dB, 1 );
magma_dsetvector( maxn, C, 1, dC1, 1 );
magma_dsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == 'N' ? m : n);
cublasDaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMV( m, n ) / 1e9;
printf( "dgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DSYMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetvector( m, B, 1, dB, 1 );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMV( m ) / 1e9;
printf( "dsymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_dlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_dgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_dsetmatrix( m, m, LU, ld, dA, ld );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "dtrsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test DGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == 'N');
bool ntb = (trans[ib] == 'N');
magma_dsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_dsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMM( m, n, k ) / 1e9;
printf( "dgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetmatrix( m, n, B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMM( side[is], m, n ) / 1e9;
printf( "dsymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYRK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_dsetmatrix( n, k, A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYRK( k, n ) / 1e9;
printf( "dsyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYR2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == 'N');
magma_dsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYR2K( k, n ) / 1e9;
printf( "dsyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRMM( side[is], m, n ) / 1e9;
printf( "dtrmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test DTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( side[is], m, n ) / 1e9;
printf( "dtrsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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;
}
