// On input, A and ipiv is LU factorization of A. On output, A is overwritten.
// Requires m == n.
// Uses init_matrix() to re-generate original A as needed.
// Generates random RHS b and solves Ax=b.
// Returns residual, |Ax - b| / (n |A| |x|).
double get_residual(
magma_int_t m, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
magma_int_t *ipiv )
{
if ( m != n ) {
printf( "\nERROR: residual check defined only for square matrices\n" );
return -1;
}
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magma_int_t ione = 1;
// this seed should be DIFFERENT than used in init_matrix
// (else x is column of A, so residual can be exactly zero)
magma_int_t ISEED[4] = {0,0,0,2};
magma_int_t info = 0;
magmaDoubleComplex *x, *b;
// initialize RHS
TESTING_MALLOC( x, magmaDoubleComplex, n );
TESTING_MALLOC( b, magmaDoubleComplex, n );
lapackf77_zlarnv( &ione, ISEED, &n, b );
blasf77_zcopy( &n, b, &ione, x, &ione );
// solve Ax = b
lapackf77_zgetrs( "Notrans", &n, &ione, A, &lda, ipiv, x, &n, &info );
if (info != 0)
printf("lapackf77_zgetrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// reset to original A
init_matrix( m, n, A, lda );
// compute r = Ax - b, saved in b
blasf77_zgemv( "Notrans", &m, &n, &c_one, A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
double norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_zlange( "F", &m, &n, A, &lda, work );
norm_r = lapackf77_zlange( "F", &n, &ione, b, &n, work );
norm_x = lapackf77_zlange( "F", &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_zprint( 1, n, b, 1 );
TESTING_FREE( x );
TESTING_FREE( b );
//printf( "r=%.2e, A=%.2e, x=%.2e, n=%d\n", norm_r, norm_A, norm_x, n );
return norm_r / (n * norm_A * norm_x);
}
// On input, LU and ipiv is LU factorization of A. On output, LU is overwritten.
// Works for any m, n.
// Uses init_matrix() to re-generate original A as needed.
// Returns error in factorization, |PA - LU| / (n |A|)
// This allocates 3 more matrices to store A, L, and U.
double get_LU_error(magma_int_t M, magma_int_t N,
magmaDoubleComplex *LU, magma_int_t lda,
magma_int_t *ipiv)
{
magma_int_t min_mn = min(M,N);
magma_int_t ione = 1;
magma_int_t i, j;
magmaDoubleComplex alpha = MAGMA_Z_ONE;
magmaDoubleComplex beta = MAGMA_Z_ZERO;
magmaDoubleComplex *A, *L, *U;
double work[1], matnorm, residual;
TESTING_MALLOC( A, magmaDoubleComplex, lda*N );
TESTING_MALLOC( L, magmaDoubleComplex, M*min_mn );
TESTING_MALLOC( U, magmaDoubleComplex, min_mn*N );
memset( L, 0, M*min_mn*sizeof(magmaDoubleComplex) );
memset( U, 0, min_mn*N*sizeof(magmaDoubleComplex) );
// set to original A
init_matrix( M, N, A, lda );
lapackf77_zlaswp( &N, A, &lda, &ione, &min_mn, ipiv, &ione);
// copy LU to L and U, and set diagonal to 1
lapackf77_zlacpy( MagmaLowerStr, &M, &min_mn, LU, &lda, L, &M );
lapackf77_zlacpy( MagmaUpperStr, &min_mn, &N, LU, &lda, U, &min_mn );
for(j=0; j<min_mn; j++)
L[j+j*M] = MAGMA_Z_MAKE( 1., 0. );
matnorm = lapackf77_zlange("f", &M, &N, A, &lda, work);
blasf77_zgemm("N", "N", &M, &N, &min_mn,
&alpha, L, &M, U, &min_mn, &beta, LU, &lda);
for( j = 0; j < N; j++ ) {
for( i = 0; i < M; i++ ) {
LU[i+j*lda] = MAGMA_Z_SUB( LU[i+j*lda], A[i+j*lda] );
}
}
residual = lapackf77_zlange("f", &M, &N, LU, &lda, work);
TESTING_FREE(A);
TESTING_FREE(L);
TESTING_FREE(U);
return residual / (matnorm * N);
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time /*cpu_time*/;
float *h_A, *h_R, *h_B, *h_S, *h_work;
float *w1, *w2, vl=0, vu=0;
float result[2] = {0};
magma_int_t *iwork;
magma_int_t N, n2, info, il, iu, m1, m2, nb, lwork, liwork;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork;
#endif
//float d_one = 1.;
//float d_ten = 10.;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf(" N M GPU Time (sec)\n");
printf("============================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
nb = magma_get_ssytrd_nb(N);
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = 2*N*nb + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = 1 + 6*N*nb + 2* N*N;
#endif
liwork = 3 + 5*N;
if ( opts.fraction == 0 ) {
il = N / 10;
iu = N / 5+il;
}
else {
il = 1;
iu = (int) (opts.fraction*N);
if (iu < 1) iu = 1;
}
TESTING_MALLOC( h_A, float, n2 );
TESTING_MALLOC( h_B, float, n2 );
TESTING_MALLOC( w1, float, N );
TESTING_MALLOC( w2, float, N );
TESTING_MALLOC( iwork, magma_int_t, liwork );
TESTING_HOSTALLOC( h_R, float, n2 );
TESTING_HOSTALLOC( h_S, float, n2 );
TESTING_HOSTALLOC( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTALLOC( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
for(int i=0; i<N; i++) {
MAGMA_S_SET2REAL( h_B[i*N+i], ( MAGMA_S_REAL(h_B[i*N+i]) + 1.*N ) );
MAGMA_S_SET2REAL( h_A[i*N+i], MAGMA_S_REAL(h_A[i*N+i]) );
}
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if(opts.warmup){
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_ssygvdx( opts.itype, opts.jobz, 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_ssygvdx( opts.itype, opts.jobz, 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | S(with V) - S(w/o V) | / | S |
=================================================================== */
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = h_work + N*N;
#endif
float temp1, temp2;
result[0] = 1.;
result[0] /= lapackf77_slansy("1", &opts.uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_slange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_ssygvdx( opts.itype, 'N', 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for(int j=0; j < m2; j++) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[1] = temp2 / (((float)m2)*temp1);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %7.2f\n",
(int) N, (int) m1, gpu_time);
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
else if (opts.itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
else if (opts.itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
printf( "(2) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e%s\n\n", result[1], (result[1] < tolulp ? "" : " failed"));
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( w1 );
TESTING_FREE( w2 );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTFREE( rwork);
#endif
TESTING_FREE( iwork );
TESTING_HOSTFREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_HOSTFREE( h_S );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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 *d_A;
float *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_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) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC( h_A, magmaFloatComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaFloatComplex, n2 );
TESTING_DEVALLOC( d_A, magmaFloatComplex, ldda*N );
TESTING_DEVALLOC( dtau, magmaFloatComplex, min_mn );
TESTING_DEVALLOC(dwork, float, min_mn );
TESTING_MALLOC( 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 );
// 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_cgeqrf 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\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, error );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( dtau );
TESTING_DEVFREE( dwork );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
int main( int argc, char** argv)
{
#define hA(i,j) (hA + (i) + (j)*lda)
TESTING_CUDA_INIT();
cuDoubleComplex c_zero = MAGMA_Z_ZERO;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex *hA, *hR, *dA;
//real_Double_t gpu_time, gpu_perf;
//int ione = 1;
//int ISEED[4] = {0, 0, 0, 1};
int nsize[] = { 32, 64, 96, 256, 100, 200, 512 };
int ntest = sizeof(nsize) / sizeof(int);
int n = nsize[ntest-1];
int lda = ((n + 31)/32)*32;
int ntile, nb;
TESTING_MALLOC ( hA, cuDoubleComplex, lda*n );
TESTING_MALLOC ( hR, cuDoubleComplex, lda*n );
TESTING_DEVALLOC ( dA, cuDoubleComplex, lda*n );
for( int t = 0; t < ntest; ++t ) {
n = nsize[t];
lda = ((n + 31)/32)*32;
// initialize matrices; entries are (i.j) for A
double nf = 100.;
for( int j = 0; j < n; ++j ) {
// upper
for( int i = 0; i < j; ++i ) {
*hA(i,j) = MAGMA_Z_MAKE( (i + j/nf)/nf, 0. );
}
// lower
for( int i = j; i < n; ++i ) {
*hA(i,j) = MAGMA_Z_MAKE( i + j/nf, 0. );
}
}
printf( "A%d = ", n );
magma_zprint( n, n, hA, lda );
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize( MagmaLower, n, dA, lda );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d = ", n );
magma_zprint( n, n, hR, lda );
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize( MagmaUpper, n, dA, lda );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "U%d = ", n );
magma_zprint( n, n, hR, lda );
// -----
//lapackf77_zlaset( "u", &n, &n, &c_zero, &c_one, hA, &lda );
nb = 64;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 32;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
ntile = (n - nb < 0 ? 0 : (n - nb) / (2*nb) + 1);
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, 2*nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d_2m = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 25;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 25;
ntile = (n - nb < 0 ? 0 : (n - nb) / (3*nb) + 1);
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, 3*nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d_3n = ", n, nb );
magma_zprint( n, n, hR, lda );
nb = 100;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaLower, nb, dA, lda, ntile, nb, nb );
magmablas_zsymmetrize( MagmaLower, n%nb, &dA[ ntile*nb*(1+lda) ], lda ); // last partial block
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "L%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
// -----
nb = 64;
ntile = n / nb;
magma_zsetmatrix( n, n, hA, lda, dA, lda );
magmablas_zsymmetrize_tiles( MagmaUpper, nb, dA, lda, ntile, nb, nb );
magma_zgetmatrix( n, n, dA, lda, hR, lda );
printf( "U%d_%d = ", n, nb );
magma_zprint( n, n, hR, lda );
}
TESTING_FREE( hA );
TESTING_FREE( hR );
TESTING_DEVFREE( dA );
/* Shutdown */
TESTING_CUDA_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cherk
*/
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 N, K;
magma_int_t Ak, An;
magma_int_t sizeA, sizeC;
magma_int_t lda, ldc, ldda, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaFloatComplex *h_A, *h_C, *h_Ccublas;
magmaFloatComplex *d_A, *d_C;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float alpha = MAGMA_D_MAKE( 0.29, -0.86 );
float beta = MAGMA_D_MAKE( -0.48, 0.38 );
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"uplo = %c, transA = %c\n", opts.uplo, opts.transA );
printf(" N K CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("==================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
K = opts.ksize[i];
gflops = FLOPS_CHERK(K, N) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = An = N;
Ak = K;
} else {
lda = An = K;
Ak = N;
}
ldc = N;
ldda = ((lda+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*Ak;
sizeC = ldc*N;
TESTING_MALLOC( h_A, magmaFloatComplex, lda*Ak );
TESTING_MALLOC( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC( h_Ccublas, magmaFloatComplex, ldc*N );
TESTING_DEVALLOC( d_A, magmaFloatComplex, ldda*Ak );
TESTING_DEVALLOC( d_C, magmaFloatComplex, lddc*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_csetmatrix( An, Ak, h_A, lda, d_A, ldda );
magma_csetmatrix( N, N, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasCherk( opts.uplo, opts.transA, N, K,
alpha, d_A, ldda,
beta, d_C, lddc );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetmatrix( N, N, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_cherk( &opts.uplo, &opts.transA, &N, &K,
&alpha, h_A, &lda,
&beta, h_C, &ldc );
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_clanhe("fro", &opts.uplo, &N, h_C, &ldc, work);
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Ccublas, &ione );
cublas_error = lapackf77_clanhe( "fro", &opts.uplo, &N, h_Ccublas, &ldc, work ) / Cnorm;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) N, (int) K,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error );
}
else {
printf("%5d %5d %7.2f (%7.2f) --- ( --- ) --- ---\n",
(int) N, (int) K,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE( h_A );
TESTING_FREE( h_C );
TESTING_FREE( h_Ccublas );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_C );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqlf
*/
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];
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_zgeqrf_nb(M);
gflops = FLOPS_ZGEQLF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgeqlf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, N*nb );
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC( h_work, magmaDoubleComplex, lwork );
/* 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
=================================================================== */
gpu_time = magma_wtime();
magma_zgeqlf( M, N, h_R, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgeqlf(&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_zgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
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\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_FREE( h_work );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrmv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
double cublas_error, Cnorm, work[1];
magma_int_t N;
magma_int_t Ak;
magma_int_t sizeA;
magma_int_t lda, ldda;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex *h_A, *h_x, *h_xcublas;
magmaDoubleComplex *d_A, *d_x;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"uplo = %c, transA = %c, diag = %c \n", opts.uplo, opts.transA, opts.diag );
printf(" N CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("==================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
gflops = FLOPS_ZTRMM(opts.side, N, 1) / 1e9;
lda = N;
Ak = N;
ldda = ((lda+31)/32)*32;
sizeA = lda*Ak;
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( h_x, magmaDoubleComplex, N );
TESTING_MALLOC( h_xcublas, magmaDoubleComplex, N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_DEVALLOC( d_x, magmaDoubleComplex, N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &N, h_x );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_zsetvector( N, h_x, 1, d_x, 1 );
cublas_time = magma_sync_wtime( NULL );
cublasZtrmv( opts.uplo, opts.transA, opts.diag,
N,
d_A, ldda,
d_x, 1 );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetvector( N, d_x, 1, h_xcublas, 1 );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ztrmv( &opts.uplo, &opts.transA, &opts.diag,
&N,
h_A, &lda,
h_x, &ione );
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_zlange( "M", &N, &ione, h_x, &N, work );
blasf77_zaxpy( &N, &c_neg_one, h_x, &ione, h_xcublas, &ione );
cublas_error = lapackf77_zlange( "M", &N, &ione, h_xcublas, &N, work ) / Cnorm;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) N,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error );
}
else {
printf("%5d %7.2f (%7.2f) --- ( --- ) --- ---\n",
(int) N,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE( h_A );
TESTING_FREE( h_x );
TESTING_FREE( h_xcublas );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_x );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ctrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
float cublas_error, normA, normx, normr, work[1];
magma_int_t N, info;
magma_int_t sizeA;
magma_int_t lda, ldda;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *ipiv;
magmaFloatComplex *h_A, *h_b, *h_x, *h_xcublas;
magmaFloatComplex *d_A, *d_x;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("uplo = %c, transA = %c, diag = %c\n", opts.uplo, opts.transA, opts.diag );
printf(" N CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("============================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
gflops = FLOPS_CTRSM(opts.side, N, 1) / 1e9;
lda = N;
ldda = ((lda+31)/32)*32;
sizeA = lda*N;
TESTING_MALLOC( ipiv, magma_int_t, N );
TESTING_MALLOC( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC( h_b, magmaFloatComplex, N );
TESTING_MALLOC( h_x, magmaFloatComplex, N );
TESTING_MALLOC( h_xcublas, magmaFloatComplex, N );
TESTING_DEVALLOC( d_A, magmaFloatComplex, ldda*N );
TESTING_DEVALLOC( d_x, magmaFloatComplex, N );
/* Initialize the matrices */
/* Factor A into LU to get well-conditioned triangular matrix.
* Copy L to U, since L seems okay when used with non-unit diagonal
* (i.e., from U), while U fails when used with unit diagonal. */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_cgetrf( &N, &N, h_A, &lda, ipiv, &info );
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < j; ++i ) {
*h_A(i,j) = *h_A(j,i);
}
}
lapackf77_clarnv( &ione, ISEED, &N, h_b );
blasf77_ccopy( &N, h_b, &ione, h_x, &ione );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
magma_csetvector( N, h_x, 1, d_x, 1 );
cublas_time = magma_sync_wtime( NULL );
cublasCtrsv( opts.uplo, opts.transA, opts.diag,
N,
d_A, ldda,
d_x, 1 );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetvector( N, d_x, 1, h_xcublas, 1 );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ctrsv( &opts.uplo, &opts.transA, &opts.diag,
&N,
h_A, &lda,
h_x, &ione );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
// error for CUBLAS
normA = lapackf77_clange( "F", &N, &N, h_A, &lda, work );
normx = lapackf77_clange( "F", &N, &ione, h_xcublas, &ione, work );
blasf77_ctrmv( &opts.uplo, &opts.transA, &opts.diag,
&N,
h_A, &lda,
h_xcublas, &ione );
blasf77_caxpy( &N, &c_neg_one, h_b, &ione, h_xcublas, &ione );
normr = lapackf77_clange( "F", &N, &ione, h_xcublas, &N, work );
cublas_error = normr / (normA*normx);
if ( opts.lapack ) {
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) N,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error );
}
else {
printf("%5d %7.2f (%7.2f) --- ( --- ) %8.2e\n",
(int) N,
cublas_perf, 1000.*cublas_time,
cublas_error );
}
TESTING_FREE( h_A );
TESTING_FREE( h_x );
TESTING_FREE( h_xcublas );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_x );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgemm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
double magma_error, cublas_error, Cnorm, work[1];
magma_int_t M, N, K;
magma_int_t Am, An, Bm, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex *h_A, *h_B, *h_C, *h_Cmagma, *h_Ccublas;
magmaDoubleComplex *d_A, *d_B, *d_C;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.48, 0.38 );
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"transA = %c, transB = %c\n", opts.transA, opts.transB );
printf(" M N K MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("=========================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
K = opts.ksize[i];
gflops = FLOPS_ZGEMM( M, N, K ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if ( opts.transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
ldc = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*An;
sizeB = ldb*Bn;
sizeC = ldc*N;
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*An );
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*Bn );
TESTING_MALLOC( h_C, magmaDoubleComplex, ldc*N );
TESTING_MALLOC( h_Cmagma, magmaDoubleComplex, ldc*N );
TESTING_MALLOC( h_Ccublas, magmaDoubleComplex, ldc*N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*An );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*Bn );
TESTING_DEVALLOC( d_C, magmaDoubleComplex, lddc*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_zlarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_zsetmatrix( Am, An, h_A, lda, d_A, ldda );
magma_zsetmatrix( Bm, Bn, h_B, ldb, d_B, lddb );
magma_zsetmatrix( M, N, h_C, ldc, d_C, lddc );
magma_time = magma_sync_wtime( NULL );
magmablas_zgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, ldda,
d_B, lddb,
beta, d_C, lddc );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, N, d_C, lddc, h_Cmagma, ldc );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( M, N, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasZgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, ldda,
d_B, lddb,
beta, d_C, lddc );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( M, N, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_zgemm( &opts.transA, &opts.transB, &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
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_zlange( "M", &M, &N, h_C, &ldc, work );
blasf77_zaxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cmagma, &ione );
magma_error = lapackf77_zlange( "M", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
blasf77_zaxpy( &sizeC, &c_neg_one, h_C, &ione, h_Ccublas, &ione );
cublas_error = lapackf77_zlange( "M", &M, &N, h_Ccublas, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
}
else {
// compute relative error for magma, relative to cublas
Cnorm = lapackf77_zlange( "M", &M, &N, h_Ccublas, &ldc, work );
blasf77_zaxpy( &sizeC, &c_neg_one, h_Ccublas, &ione, h_Cmagma, &ione );
magma_error = lapackf77_zlange( "M", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e ---\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error );
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( h_C );
TESTING_FREE( h_Cmagma );
TESTING_FREE( h_Ccublas );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
TESTING_DEVFREE( d_C );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgelqf
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
float flops, gpu_perf, cpu_perf;
float matnorm, work[1];
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
cuFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
/* Matrix size */
magma_int_t M = 0, N = 0, n2, lda, lwork;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (N>0 && M>0)
printf(" testing_cgelqf -M %d -N %d\n\n", (int) M, (int) N);
else
{
printf("\nUsage: \n");
printf(" testing_cgelqf -M %d -N %d\n\n", (int) M, (int) N);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_cgelqf -M %d -N %d\n\n", 1024, 1024);
M = N = size[9];
}
n2 = M * N;
min_mn = min(M, N);
nb = magma_get_cgeqrf_nb(M);
TESTING_MALLOC( tau, cuFloatComplex, min_mn );
TESTING_MALLOC( h_A, cuFloatComplex, n2 );
TESTING_HOSTALLOC( h_R, cuFloatComplex, n2 );
lwork = -1;
lapackf77_cgelqf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max( lwork, M*nb );
TESTING_HOSTALLOC( h_work, cuFloatComplex, lwork );
printf(" M N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("==========================================================\n");
for(i=0; i<10; i++){
if (argc == 1){
M = N = size[i];
}
min_mn= min(M, N);
lda = M;
n2 = lda*N;
flops = FLOPS( (float)M, (float)N ) / 1000000;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_cgelqf( M, N, h_R, lda, tau, h_work, lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of magma_cgelqf had an illegal value.\n", (int) -info);
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_cgelqf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of lapack_cgelqf had an illegal value.\n", (int) -info);
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
matnorm = lapackf77_clange("f", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %5d %6.2f %6.2f %e\n",
(int) M, (int) N, cpu_perf, gpu_perf,
lapackf77_clange("f", &M, &N, h_R, &lda, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_HOSTFREE( h_work );
/* Shutdown */
TESTING_CUDA_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 *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_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" 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];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_ZPOTRF( N ) / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_DEVALLOC( 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 );
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zpotf2_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - 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( &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 );
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\n",
(int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., error );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing clarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_CUDA_INIT();
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
printf( "\nUsage: %s -M m -N n -K k\n\n", argv[0] );
magma_int_t m = 500;
magma_int_t n = 300;
magma_int_t k = 32;
for( int i = 1; i < argc; i++ ) {
if (strcmp("-M", argv[i]) == 0 && i+1 < argc) {
m = atoi( argv[++i] );
}
else if (strcmp("-N", argv[i]) == 0 && i+1 < argc) {
n = atoi( argv[++i] );
}
else if (strcmp("-K", argv[i]) == 0 && i+1 < argc) {
k = atoi( argv[++i] );
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( k <= 0 || k > m || k > n ) {
printf( "requires 0 < k <= min(m,n)\n" );
exit(1);
}
magma_int_t ldc = m;
magma_int_t ldv = max(m,n);
magma_int_t ldt = k;
magma_int_t ldw = max(m,n);
magma_int_t nv;
ldc = ((ldc+31)/32)*32;
ldv = ((ldv+31)/32)*32;
ldt = ((ldt+31)/32)*32;
ldw = ((ldw+31)/32)*32;
// Allocate memory for matrices
cuFloatComplex *C, *R, *V, *T, *W;
TESTING_MALLOC( C, cuFloatComplex, ldc*n );
TESTING_MALLOC( R, cuFloatComplex, ldc*n );
TESTING_MALLOC( V, cuFloatComplex, ldv*k );
TESTING_MALLOC( T, cuFloatComplex, ldt*k );
TESTING_MALLOC( W, cuFloatComplex, ldw*k );
cuFloatComplex *dC, *dV, *dT, *dW;
TESTING_DEVALLOC( dC, cuFloatComplex, ldc*n );
TESTING_DEVALLOC( dV, cuFloatComplex, ldv*k );
TESTING_DEVALLOC( dT, cuFloatComplex, ldt*k );
TESTING_DEVALLOC( dW, cuFloatComplex, ldw*k );
magma_int_t size;
magma_int_t iseed[4] = { 1, 2, 3, 4 };
float error, work[1];
// test all combinations of input parameters
const char* side[] = { MagmaLeftStr, MagmaRightStr };
const char* trans[] = { MagmaConjTransStr, MagmaNoTransStr };
const char* direct[] = { MagmaForwardStr, MagmaBackwardStr };
const char* storev[] = { MagmaColumnwiseStr, MagmaRowwiseStr };
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("==================================================================================\n");
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
//printf( "# ----------\n" );
//printf( "# %-10s %-10s %-10s %-10s\n", storev[istor], side[iside], direct[idir], trans[itran] );
// C is full
size = ldc*n;
lapackf77_clarnv( &ione, iseed, &size, C );
//printf( "C=" ); magma_cprint( m, n, C, ldc );
// V is ldv x nv. See larfb docs for description.
ldv = (*side[iside] == 'L' ? m : n);
nv = k;
size = ldv*nv;
lapackf77_clarnv( &ione, iseed, &size, V );
if ( *storev[istor] == MagmaColumnwise ) {
if ( *direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaUpperStr, &k, &k, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaLowerStr, &k, &k, &c_zero, &c_one, &V[(ldv-k)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( *direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &k, &k, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_claset( MagmaUpperStr, &k, &k, &c_zero, &c_one, &V[(nv-k)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_cprint( ldv, nv, V, ldv );
// T is upper triangular for forward, and lower triangular for backward
magma_int_t k1 = k-1;
size = ldt*k;
lapackf77_clarnv( &ione, iseed, &size, T );
if ( *direct[idir] == MagmaForward ) {
lapackf77_claset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_claset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_cprint( k, k, T, ldt );
magma_csetmatrix( m, n, C, ldc, dC, ldc );
magma_csetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_csetmatrix( k, k, T, ldt, dT, ldt );
lapackf77_clarfb( side[iside], trans[itran], direct[idir], storev[istor],
&m, &n, &k,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_cprint( m, n, C, ldc );
magma_clarfb_gpu( *side[iside], *trans[itran], *direct[idir], *storev[istor],
m, n, k,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_cgetmatrix( m, n, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_cprint( m, n, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %-10s %-10s %-10s %-10s %8.2e\n",
(int) m, (int) n, (int) k,
storev[istor], side[iside], direct[idir], trans[itran], error );
}}}}
// Memory clean up
TESTING_FREE( C );
TESTING_FREE( R );
TESTING_FREE( V );
TESTING_FREE( T );
TESTING_FREE( W );
TESTING_DEVFREE( dC );
TESTING_DEVFREE( dV );
TESTING_DEVFREE( dT );
TESTING_DEVFREE( dW );
// Shutdown
TESTING_CUDA_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double magma_error, cublas_error, work[1];
magma_int_t M, N, info;
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};
magma_int_t *piv;
magma_err_t err;
magmaDoubleComplex *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2, *LU, *LUT;
magmaDoubleComplex *d_A, *d_B;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"side = %c, uplo = %c, transA = %c, diag = %c \n", opts.side, opts.uplo, opts.transA, opts.diag );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("==================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
gflops = FLOPS_ZTRSM(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( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( LU, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( LUT, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_B1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_X1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_X2, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_Bcublas, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_Bmagma, magmaDoubleComplex, ldb*N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, LU );
err = magma_malloc_cpu( (void**) &piv, Ak*sizeof(magma_int_t) ); assert( err == 0 );
lapackf77_zgetrf( &Ak, &Ak, LU, &lda, piv, &info );
int i, j;
for(i=0;i<Ak;i++){
for(j=0;j<Ak;j++){
LUT[j+i*lda] = LU[i+j*lda];
}
}
lapackf77_zlacpy(MagmaUpperStr, &Ak, &Ak, LUT, &lda, LU, &lda);
if(opts.uplo == MagmaLower){
lapackf77_zlacpy(MagmaLowerStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}else{
lapackf77_zlacpy(MagmaUpperStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(magmaDoubleComplex));
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_zsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_ztrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasZtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ztrsm( &opts.side, &opts.uplo, &opts.transA, &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
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(magmaDoubleComplex));
magmaDoubleComplex alpha2 = MAGMA_Z_DIV( c_one, alpha );
blasf77_ztrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
double norm1 = lapackf77_zlange( "M", &M, &N, h_X1, &ldb, work );
double normx = lapackf77_zlange( "M", &M, &N, h_Bmagma, &ldb, work );
double normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(magmaDoubleComplex));
blasf77_ztrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_zlange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_zlange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error );
}
TESTING_FREE( h_A );
TESTING_FREE( LU );
TESTING_FREE( LUT );
TESTING_FREE( h_B );
TESTING_FREE( h_Bcublas );
TESTING_FREE( h_Bmagma );
TESTING_FREE( h_B1 );
TESTING_FREE( h_X1 );
TESTING_FREE( h_X2 );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeadd_batched
Code is very similar to testing_clacpy_batched.cpp
*/
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_B;
magmaFloatComplex *d_A, *d_B;
magmaFloatComplex **hAarray, **hBarray, **dAarray, **dBarray;
magmaFloatComplex alpha = MAGMA_C_MAKE( 3.1415, 2.718 );
magma_int_t M, N, mb, nb, size, lda, ldda, mstride, nstride, ntile;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
mb = (opts.nb == 0 ? 32 : opts.nb);
nb = (opts.nb == 0 ? 64 : opts.nb);
mstride = 2*mb;
nstride = 3*nb;
printf("mb=%d, nb=%d, mstride=%d, nstride=%d\n", (int) mb, (int) nb, (int) mstride, (int) nstride );
printf(" M N ntile CPU GFlop/s (sec) GPU GFlop/s (sec) error \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];
lda = M;
ldda = ((M+31)/32)*32;
size = lda*N;
if ( N < nb || M < nb ) {
ntile = 0;
} else {
ntile = min( (M - nb)/mstride + 1,
(N - nb)/nstride + 1 );
}
gflops = 2.*mb*nb*ntile / 1e9;
TESTING_MALLOC( h_A, magmaFloatComplex, lda *N );
TESTING_MALLOC( h_B, magmaFloatComplex, lda *N );
TESTING_DEVALLOC( d_A, magmaFloatComplex, ldda*N );
TESTING_DEVALLOC( d_B, magmaFloatComplex, ldda*N );
TESTING_MALLOC( hAarray, magmaFloatComplex*, ntile );
TESTING_MALLOC( hBarray, magmaFloatComplex*, ntile );
TESTING_DEVALLOC( dAarray, magmaFloatComplex*, ntile );
TESTING_DEVALLOC( dBarray, magmaFloatComplex*, ntile );
lapackf77_clarnv( &ione, ISEED, &size, h_A );
lapackf77_clarnv( &ione, ISEED, &size, h_B );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
magma_csetmatrix( M, N, h_B, lda, d_B, ldda );
// setup pointers
for( int tile = 0; tile < ntile; ++tile ) {
int offset = tile*mstride + tile*nstride*ldda;
hAarray[tile] = &d_A[offset];
hBarray[tile] = &d_B[offset];
}
magma_setvector( ntile, sizeof(magmaFloatComplex*), hAarray, 1, dAarray, 1 );
magma_setvector( ntile, sizeof(magmaFloatComplex*), hBarray, 1, dBarray, 1 );
gpu_time = magma_sync_wtime( 0 );
magmablas_cgeadd_batched( mb, nb, alpha, dAarray, ldda, dBarray, ldda, ntile );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for( int tile = 0; tile < ntile; ++tile ) {
int offset = tile*mstride + tile*nstride*lda;
for( int j = 0; j < nb; ++j ) {
blasf77_caxpy( &mb, &alpha,
&h_A[offset + j*lda], &ione,
&h_B[offset + j*lda], &ione );
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_cgetmatrix( M, N, d_B, ldda, h_A, lda );
error = lapackf77_clange( "F", &M, &N, h_B, &lda, work );
blasf77_caxpy(&size, &c_neg_one, h_A, &ione, h_B, &ione);
error = lapackf77_clange("f", &M, &N, h_B, &lda, work) / error;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, (int) ntile,
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
TESTING_FREE( hAarray );
TESTING_FREE( hBarray );
TESTING_DEVFREE( dAarray );
TESTING_DEVFREE( dBarray );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cunmqr
*/
int main( int argc, char** argv )
{
TESTING_CUDA_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
/* Matrix size */
magma_int_t m, n, k;
const int MAXTESTS = 10;
magma_int_t msize[MAXTESTS] = { 1024, 2048, 3072, 4032, 5184, 6016, 7040, 8064, 9088, 10112 };
magma_int_t nsize[MAXTESTS] = { 1024, 2048, 3072, 4032, 5184, 6016, 7040, 8064, 9088, 10112 };
magma_int_t ksize[MAXTESTS] = { 1024, 2048, 3072, 4032, 5184, 6016, 7040, 8064, 9088, 10112 };
magma_int_t size;
magma_int_t info;
magma_int_t iseed[4] = {0,0,0,1};
printf( "Usage: %s -N m,n,k -c\n"
" -N can be repeated %d times. m > 0, n > 0, k > 0 is required.\n"
" If only m,n is given, then n=k. If only m is given, then m=n=k.\n"
" -c or setting $MAGMA_TESTINGS_CHECK runs LAPACK and checks result.\n\n",
argv[0], MAXTESTS );
int checkres = (getenv("MAGMA_TESTINGS_CHECK") != NULL);
int ntest = 0;
magma_int_t nmax = 0;
magma_int_t mmax = 0;
magma_int_t kmax = 0;
for( int i = 1; i < argc; i++ ) {
if ( strcmp("-N", argv[i]) == 0 && i+1 < argc ) {
magma_assert( ntest < MAXTESTS, "error: -N repeated more than maximum %d tests\n", MAXTESTS );
info = sscanf( argv[++i], "%d,%d,%d", &m, &n, &k );
if ( info == 3 && m > 0 && n > 0 && k > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = n;
ksize[ ntest ] = k;
}
else if ( info == 2 && m > 0 && n > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = n;
ksize[ ntest ] = n; // implicitly
}
else if ( info == 1 && m > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = m; // implicitly
ksize[ ntest ] = m; // implicitly
}
else {
printf( "error: -N %s is invalid; ensure m > 0, n > 0, k > 0.\n", argv[i] );
exit(1);
}
mmax = max( mmax, msize[ntest] );
nmax = max( nmax, nsize[ntest] );
kmax = max( kmax, ksize[ntest] );
ntest++;
}
else if ( strcmp("-c", argv[i]) == 0 ) {
checkres = true;
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( ntest == 0 ) {
ntest = MAXTESTS;
nmax = nsize[ntest-1];
mmax = msize[ntest-1];
kmax = ksize[ntest-1];
}
m = mmax;
n = nmax;
k = kmax;
assert( n > 0 && m > 0 && k > 0 );
magma_int_t nb = magma_get_cgeqrf_nb( m );
magma_int_t ldc = m;
magma_int_t lda = max(m,n);
ldc = ((ldc+31)/32)*32;
lda = ((lda+31)/32)*32;
// Allocate memory for matrices
cuFloatComplex *C, *R, *A, *W, *tau;
magma_int_t lwork = max( m*nb, n*nb );
magma_int_t lwork_max = lwork;
TESTING_MALLOC( C, cuFloatComplex, ldc*n );
TESTING_MALLOC( R, cuFloatComplex, ldc*n );
TESTING_MALLOC( A, cuFloatComplex, lda*k );
TESTING_MALLOC( W, cuFloatComplex, lwork_max );
TESTING_MALLOC( tau, cuFloatComplex, k );
// test all combinations of input parameters
const char* side[] = { MagmaLeftStr, MagmaRightStr };
const char* trans[] = { MagmaConjTransStr, MagmaNoTransStr };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int i = 0; i < ntest; ++i ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
m = msize[i];
n = nsize[i];
k = ksize[i];
if ( *side[iside] == 'L' && m < k ) {
printf( "%5d %5d %5d %-5s %-9s skipping because side=left and m < k\n",
(int) m, (int) n, (int) k, side[iside], trans[itran] );
continue;
}
if ( *side[iside] == 'R' && n < k ) {
printf( "%5d %5d %5d %-5s %-9s skipping because side=right and n < k\n",
(int) m, (int) n, (int) k, side[iside], trans[itran] );
continue;
}
gflops = FLOPS_CUNMQR( m, n, k, *side[iside] ) / 1e9;
// C is full, m x n
size = ldc*n;
lapackf77_clarnv( &ione, iseed, &size, C );
lapackf77_clacpy( "Full", &m, &n, C, &ldc, R, &ldc );
//magma_csetmatrix( m, n, C, ldc, dC, ldc );
// A is m x k (left) or n x k (right)
lda = (*side[iside] == 'L' ? m : n);
size = lda*k;
lapackf77_clarnv( &ione, iseed, &size, A );
// compute QR factorization to get Householder vectors in A, tau
magma_cgeqrf( lda, k, A, lda, tau, W, lwork_max, &info );
if ( info != 0 )
printf("magma_cgeqrf returned error %d\n", info);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cunmqr( side[iside], trans[itran],
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cunmqr returned error %d.\n", (int) info);
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for work size
lwork = -1;
magma_cunmqr( *side[iside], *trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
if (info != 0)
printf("magma_cunmqr returned error %d (lwork query).\n", (int) info);
lwork = (magma_int_t) MAGMA_C_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max )
printf("invalid lwork %d, lwork_max %d\n", lwork, lwork_max );
gpu_time = magma_wtime();
magma_cunmqr( *side[iside], *trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cunmqr returned error %d.\n", (int) info);
//magma_cgetmatrix( m, n, dC, ldc, R, ldc );
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %-5s %-9s %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) m, (int) n, (int) k, side[iside], trans[itran],
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
}} // end iside, itran
printf( "\n" );
} // end i
// Memory clean up
TESTING_FREE( C );
TESTING_FREE( R );
TESTING_FREE( A );
TESTING_FREE( W );
TESTING_FREE( tau );
// Shutdown
TESTING_CUDA_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cheevd
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
cuFloatComplex *h_A, *h_R, *h_work;
float *rwork, *w1, *w2;
magma_int_t *iwork;
float gpu_time, cpu_time;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t size[8] = {1024,2048,3072,4032,5184,6016,7040,8064};
magma_int_t i, info;
magma_int_t ione = 1, izero = 0;
magma_int_t ISEED[4] = {0,0,0,1};
const char *uplo = MagmaLowerStr;
const char *jobz = MagmaVectorsStr;
magma_int_t checkres;
float result[3], eps = lapackf77_slamch( "E" );
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0) {
N = atoi(argv[++i]);
}
else if ( strcmp("-JV", argv[i]) == 0 ) {
jobz = MagmaVectorsStr;
}
else if ( strcmp("-JN", argv[i]) == 0 ) {
jobz = MagmaNoVectorsStr;
}
}
if (N>0)
printf(" testing_cheevd -N %d [-JV] [-JN]\n\n", (int) N);
else {
printf("\nUsage: \n");
printf(" testing_cheevd -N %d [-JV] [-JN]\n\n", (int) N);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_cheevd -N %d [-JV] [-JN]\n\n", 1024);
N = size[7];
}
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
if ( checkres && jobz[0] == MagmaNoVectors ) {
printf( "Cannot check results when vectors are not computed (jobz='N')\n" );
checkres = false;
}
/* Query for workspace sizes */
cuFloatComplex aux_work[1];
float aux_rwork[1];
magma_int_t aux_iwork[1];
magma_cheevd( jobz[0], uplo[0],
N, h_R, N, w1,
aux_work, -1,
aux_rwork, -1,
aux_iwork, -1,
&info );
magma_int_t lwork, lrwork, liwork;
lwork = (magma_int_t) MAGMA_C_REAL( aux_work[0] );
lrwork = (magma_int_t) aux_rwork[0];
liwork = aux_iwork[0];
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, cuFloatComplex, N*N );
TESTING_MALLOC( w1, float , N );
TESTING_MALLOC( w2, float , N );
TESTING_HOSTALLOC( h_R, cuFloatComplex, N*N );
TESTING_HOSTALLOC( h_work, cuFloatComplex, lwork );
TESTING_MALLOC( rwork, float, lrwork );
TESTING_MALLOC( iwork, magma_int_t, liwork );
printf(" N CPU Time(s) GPU Time(s) \n");
printf("===================================\n");
for(i=0; i<8; i++){
if (argc==1){
N = size[i];
}
n2 = N*N;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
for( int i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_A[i*N+i], MAGMA_C_REAL(h_A[i*N+i]) );
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
/* warm up run */
magma_cheevd(jobz[0], uplo[0],
N, h_R, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
/* query for optimal workspace sizes */
magma_cheevd(jobz[0], uplo[0],
N, h_R, N, w1,
h_work, -1,
rwork, -1,
iwork, -1,
&info);
int lwork_save = lwork;
int lrwork_save = lrwork;
int liwork_save = liwork;
lwork = min( lwork, (magma_int_t) MAGMA_C_REAL( h_work[0] ));
lrwork = min( lrwork, (magma_int_t) rwork[0] );
liwork = min( liwork, iwork[0] );
//printf( "lwork %d, query %d, used %d; liwork %d, query %d, used %d\n",
// lwork_save, (magma_int_t) h_work[0], lwork,
// liwork_save, iwork[0], liwork );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_cheevd(jobz[0], uplo[0],
N, h_R, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
lwork = lwork_save;
lrwork = lrwork_save;
liwork = liwork_save;
if ( checkres ) {
/* =====================================================================
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 |
=================================================================== */
float temp1, temp2;
cuFloatComplex *tau;
lapackf77_chet21(&ione, uplo, &N, &izero,
h_A, &N,
w1, w1,
h_R, &N,
h_R, &N,
tau, h_work, rwork, &result[0]);
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
magma_cheevd('N', uplo[0],
N, h_R, N, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
temp1 = temp2 = 0;
for(int j=0; j<N; j++){
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / temp1;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_cheevd(jobz, uplo,
&N, h_A, &N, w2,
h_work, &lwork,
rwork, &lrwork,
iwork, &liwork,
&info);
end = get_current_time();
if (info < 0)
printf("Argument %d of cheevd had an illegal value.\n", (int) -info);
cpu_time = GetTimerValue(start,end)/1000.;
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %6.2f %6.2f\n",
(int) N, cpu_time, gpu_time);
if ( checkres ){
printf("Testing the factorization A = U S U' for correctness:\n");
printf("(1) | A - U S U' | / (|A| N) = %e\n", result[0]*eps);
printf("(2) | I - U'U | / N = %e\n", result[1]*eps);
printf("(3) | S(w/ U)-S(w/o U)|/ |S| = %e\n\n", result[2]);
}
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( h_A);
TESTING_FREE( w1);
TESTING_FREE( w2);
TESTING_FREE( rwork);
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cpotrf
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
float flops, gpu_perf, cpu_perf;
cuFloatComplex *h_A, *h_R;
cuFloatComplex *d_A;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info;
const char *uplo = MagmaUpperStr;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float work[1], matnorm;
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_cpotri_gpu -N %d\n\n", 1024);
}
/* Allocate host memory for the matrix */
n2 = size[9] * size[9];
ldda = ((size[9]+31)/32) * 32;
TESTING_MALLOC( h_A, cuFloatComplex, n2);
TESTING_HOSTALLOC( h_R, cuFloatComplex, n2);
TESTING_DEVALLOC( d_A, cuFloatComplex, ldda*size[9] );
printf(" N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("========================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
flops = FLOPS_CPOTRI( (float)N ) / 1000000;
ldda = ((N+31)/32)*32;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_A[i*lda+i], ( MAGMA_C_REAL(h_A[i*lda+i]) + 1.*N ) );
for(j=0; j<i; j++)
h_A[i*lda+j] = cuConjf(h_A[j*lda+i]);
}
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
//cublasSetMatrix( N, N, sizeof(cuFloatComplex), h_A, lda, d_A, ldda);
//magma_cpotrf_gpu(uplo[0], N, d_A, ldda, &info);
/* factorize matrix */
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
magma_cpotrf_gpu(uplo[0], N, d_A, ldda, &info);
// check for exact singularity
//magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
//h_R[ 10 + 10*lda ] = MAGMA_C_MAKE( 0.0, 0.0 );
//magma_csetmatrix( N, N, h_R, lda, d_A, ldda );
start = get_current_time();
magma_cpotri_gpu(uplo[0], N, d_A, ldda, &info);
end = get_current_time();
if (info != 0)
printf("magma_cpotri_gpu returned error %d\n", (int) info);
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_cpotrf(uplo, &N, h_A, &lda, &info);
start = get_current_time();
lapackf77_cpotri(uplo, &N, h_A, &lda, &info);
end = get_current_time();
if (info != 0)
printf("lapackf77_cpotri returned error %d\n", (int) info);
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
matnorm = lapackf77_clange("f", &N, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %6.2f %6.2f %e\n",
(int) size[i], cpu_perf, gpu_perf,
lapackf77_clange("f", &N, &N, h_R, &lda, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeadd
*/
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 *h_A, *h_B, *d_A, *d_B;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 3.1415, 2.718 );
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t M, N, size, lda, ldda;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
/* Uncomment these lines to check parameters.
* magma_xerbla calls lapack's xerbla to print out error. */
//magmablas_zgeadd( -1, N, alpha, d_A, ldda, d_B, ldda );
//magmablas_zgeadd( M, -1, alpha, d_A, ldda, d_B, ldda );
//magmablas_zgeadd( M, N, alpha, d_A, M-1, d_B, ldda );
//magmablas_zgeadd( M, N, alpha, d_A, ldda, d_B, N-1 );
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |Bl-Bm|/|Bl|\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];
lda = M;
ldda = ((M+31)/32)*32;
size = lda*N;
gflops = 2.*M*N / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, lda *N );
TESTING_MALLOC( h_B, magmaDoubleComplex, lda *N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, ldda*N );
lapackf77_zlarnv( &ione, ISEED, &size, h_A );
lapackf77_zlarnv( &ione, ISEED, &size, h_B );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, N, h_B, lda, d_B, ldda );
gpu_time = magma_sync_wtime( NULL );
magmablas_zgeadd( M, N, alpha, d_A, ldda, d_B, ldda );
gpu_time = magma_sync_wtime( NULL ) - gpu_time;
gpu_perf = gflops / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for( int j = 0; j < N; ++j ) {
blasf77_zaxpy( &M, &alpha, &h_A[j*lda], &ione, &h_B[j*lda], &ione );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check result
=================================================================== */
magma_zgetmatrix( M, N, d_B, ldda, h_A, lda );
error = lapackf77_zlange( "F", &M, &N, h_B, &lda, work );
blasf77_zaxpy( &size, &c_neg_one, h_A, &ione, h_B, &ione );
error = lapackf77_zlange( "F", &M, &N, h_B, &lda, work ) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cunmqr
*/
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;
magma_int_t ione = 1;
magma_int_t m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
magmaFloatComplex *C, *R, *A, *W, *tau;
magma_opts opts;
parse_opts( argc, argv, &opts );
// test all combinations of input parameters
const char* side[] = { MagmaLeftStr, MagmaRightStr };
const char* trans[] = { MagmaConjTransStr, MagmaNoTransStr };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
nb = magma_get_cgeqrf_nb( m );
ldc = ((m + 31)/32)*32;
lda = ((max(m,n) + 31)/32)*32;
gflops = FLOPS_CUNMQR( m, n, k, *side[iside] ) / 1e9;
if ( *side[iside] == 'L' && m < k ) {
printf( "%5d %5d %5d %-5s %-9s skipping because side=left and m < k\n",
(int) m, (int) n, (int) k, side[iside], trans[itran] );
continue;
}
if ( *side[iside] == 'R' && n < k ) {
printf( "%5d %5d %5d %-5s %-9s skipping because side=right and n < k\n",
(int) m, (int) n, (int) k, side[iside], trans[itran] );
continue;
}
lwork_max = max( m*nb, n*nb );
TESTING_MALLOC( C, magmaFloatComplex, ldc*n );
TESTING_MALLOC( R, magmaFloatComplex, ldc*n );
TESTING_MALLOC( A, magmaFloatComplex, lda*k );
TESTING_MALLOC( W, magmaFloatComplex, lwork_max );
TESTING_MALLOC( tau, magmaFloatComplex, k );
// C is full, m x n
size = ldc*n;
lapackf77_clarnv( &ione, ISEED, &size, C );
lapackf77_clacpy( "Full", &m, &n, C, &ldc, R, &ldc );
//magma_csetmatrix( m, n, C, ldc, dC, ldc );
// A is m x k (left) or n x k (right)
lda = (*side[iside] == 'L' ? m : n);
size = lda*k;
lapackf77_clarnv( &ione, ISEED, &size, A );
// compute QR factorization to get Householder vectors in A, tau
magma_cgeqrf( lda, k, A, lda, tau, W, lwork_max, &info );
if (info != 0)
printf("magma_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cunmqr( side[iside], trans[itran],
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cunmqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for work size
lwork = -1;
magma_cunmqr( *side[iside], *trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
if (info != 0)
printf("magma_cunmqr (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_C_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max )
printf("invalid lwork %d, lwork_max %d\n", (int) lwork, (int) lwork_max );
gpu_time = magma_wtime();
magma_cunmqr( *side[iside], *trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cunmqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
//magma_cgetmatrix( m, n, dC, ldc, R, ldc );
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %-5s %-9s %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) m, (int) n, (int) k, side[iside], trans[itran],
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
TESTING_FREE( C );
TESTING_FREE( R );
TESTING_FREE( A );
TESTING_FREE( W );
TESTING_FREE( tau );
}} // end iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return 0;
}
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, cublas_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 1.5, -2.3 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.6, 0.8 );
magmaFloatComplex *A, *X, *Y, *Ycublas, *Ymagma;
magmaFloatComplex *dA, *dX, *dY, *dC_work;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("=============================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = ((N + 31)/32)*32;
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_CHEMV( N ) / 1e9;
TESTING_MALLOC( A, magmaFloatComplex, sizeA );
TESTING_MALLOC( X, magmaFloatComplex, sizeX );
TESTING_MALLOC( Y, magmaFloatComplex, sizeY );
TESTING_MALLOC( Ycublas, magmaFloatComplex, sizeY );
TESTING_MALLOC( Ymagma, magmaFloatComplex, sizeY );
TESTING_DEVALLOC( dA, magmaFloatComplex, sizeA );
TESTING_DEVALLOC( dX, magmaFloatComplex, sizeX );
TESTING_DEVALLOC( dY, magmaFloatComplex, sizeY );
blocks = (N + nb - 1) / nb;
ldwork = lda * (blocks + 1);
TESTING_DEVALLOC( dC_work, magmaFloatComplex, ldwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &sizeA, A );
magma_cmake_hermitian( N, A, lda );
lapackf77_clarnv( &ione, ISEED, &sizeX, X );
lapackf77_clarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( N, N, A, lda, dA, lda );
magma_csetvector( N, X, incx, dX, incx );
magma_csetvector( N, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasChemv( opts.uplo, N, alpha, dA, lda, dX, incx, beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetvector( N, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using MAGMA BLAS
=================================================================== */
magma_csetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
#if (GPUSHMEM >= 200)
magmablas_chemv2( opts.uplo, N, alpha, dA, lda, dX, incx, beta, dY, incy, dC_work, ldwork );
#else
magmablas_chemv( opts.uplo, N, alpha, dA, lda, dX, incx, beta, dY, incy );
#endif
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetvector( N, dY, incy, Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_chemv( &opts.uplo, &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_caxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_clange( "M", &N, &ione, Ymagma, &N, work ) / N;
blasf77_caxpy( &N, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_clange( "M", &N, &ione, Ycublas, &N, work ) / N;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
TESTING_FREE( A );
TESTING_FREE( X );
TESTING_FREE( Y );
TESTING_FREE( Ycublas );
TESTING_FREE( Ymagma );
TESTING_DEVFREE( dA );
TESTING_DEVFREE( dX );
TESTING_DEVFREE( dY );
TESTING_DEVFREE( dC_work );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magma_int_t M, N, n2, lda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol, eps = lapackf77_dlamch("E");
tol = opts.tolerance * eps;
opts.lapack |= ( opts.check == 2 ); // check (-c2) implies lapack (-l)
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("===============================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
}
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
nb = magma_get_zgeqrf_nb(M);
gflops = FLOPS_ZGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, max( N*nb, 2*nb*nb ));
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC( h_work, magmaDoubleComplex, lwork );
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgeqrf(M, N, h_R, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
magmaDoubleComplex *tau;
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_zgeqrf(&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_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE( tau );
}
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
magmaDoubleComplex *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC( h_W1, magmaDoubleComplex, n2 ); // Q
TESTING_MALLOC( h_W2, magmaDoubleComplex, n2 ); // R
TESTING_MALLOC( h_W3, magmaDoubleComplex, lwork ); // WORK
TESTING_MALLOC( h_RW, double, M ); // RWORK
lapackf77_zlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_zqrt02( &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( h_W1 );
TESTING_FREE( h_W2 );
TESTING_FREE( h_W3 );
TESTING_FREE( h_RW );
} else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
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;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf("%s\n", (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf
*/
int main( magma_int_t argc, char** argv)
{
magma_int_t nquarkthreads=2;
magma_int_t nthreads=2;
magma_int_t num_gpus = 1;
TRACE = 0;
//magma_qr_params mp;
cuDoubleComplex *h_A, *h_R, *h_work, *tau;
double gpu_perf, cpu_perf, flops;
magma_timestr_t start, end;
magma_qr_params *mp = (magma_qr_params*)malloc(sizeof(magma_qr_params));
/* Matrix size */
magma_int_t M=0, N=0, n2;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
cublasStatus status;
magma_int_t i, j, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
mp->nb=-1;
mp->ob=-1;
mp->fb=-1;
mp->ib=32;
magma_int_t loop = argc;
magma_int_t accuracyflag = 1;
char precision;
magma_int_t nc = -1;
magma_int_t ncps = -1;
if (argc != 1)
{
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
else if (strcmp("-F", argv[i])==0)
mp->fb = atoi(argv[++i]);
else if (strcmp("-O", argv[i])==0)
mp->ob = atoi(argv[++i]);
else if (strcmp("-B", argv[i])==0)
mp->nb = atoi(argv[++i]);
else if (strcmp("-b", argv[i])==0)
mp->ib = atoi(argv[++i]);
else if (strcmp("-A", argv[i])==0)
accuracyflag = atoi(argv[++i]);
else if (strcmp("-P", argv[i])==0)
nthreads = atoi(argv[++i]);
else if (strcmp("-Q", argv[i])==0)
nquarkthreads = atoi(argv[++i]);
else if (strcmp("-nc", argv[i])==0)
nc = atoi(argv[++i]);
else if (strcmp("-ncps", argv[i])==0)
ncps = atoi(argv[++i]);
}
if ((M>0 && N>0) || (M==0 && N==0))
{
printf(" testing_zgeqrf-v2 -M %d -N %d\n\n", M, N);
if (M==0 && N==0) {
M = N = size[9];
loop = 1;
}
}
else
{
printf("\nUsage: \n");
printf(" Make sure you set the number of BLAS threads to 1, e.g.,\n");
printf(" > setenv MKL_NUM_THREADS 1\n");
printf(" > testing_zgeqrf-v2 -M %d -N %d -B 128 -T 1\n\n", 1024, 1024);
exit(1);
}
}
else
{
printf("\nUsage: \n");
printf(" Make sure you set the number of BLAS threads to 1, e.g.,\n");
printf(" > setenv MKL_NUM_THREADS 1\n");
printf(" Set number of cores per socket and number of cores.\n");
printf(" > testing_zgeqrf-v2 -M %d -N %d -ncps 6 -nc 12\n\n", 1024, 1024);
printf(" Alternatively, set:\n");
printf(" Q: Number of threads for panel factorization.\n");
printf(" P: Number of threads for trailing matrix update (CPU).\n");
printf(" B: Block size.\n");
printf(" b: Inner block size.\n");
printf(" O: Block size for trailing matrix update (CPU).\n");
printf(" > testing_zgeqrf-v2 -M %d -N %d -Q 4 -P 4 -B 128 -b 32 -O 200\n\n", 10112, 10112);
M = N = size[9];
}
/* Auto tune based on number of cores and number of cores per socket if provided */
if ((nc > 0) && (ncps > 0)) {
precision = 's';
#if (defined(PRECISION_d))
precision = 'd';
#endif
#if (defined(PRECISION_c))
precision = 'c';
#endif
#if (defined(PRECISION_z))
precision = 'z';
#endif
auto_tune('q', precision, nc, ncps, M, N,
&(mp->nb), &(mp->ob), &(mp->ib), &nthreads, &nquarkthreads);
fprintf(stderr,"%d %d %d %d %d\n",mp->nb,mp->ob,mp->ib,nquarkthreads,nthreads);
}
/* Initialize MAGMA hardware context, seeting how many CPU cores
and how many GPUs to be used in the consequent computations */
mp->sync0 = 0;
magma_context *context;
context = magma_init((void*)(mp),cpu_thread, nthreads, nquarkthreads, num_gpus, argc, argv);
context->params = (void *)(mp);
mp->sync1 = (volatile magma_int_t *) malloc (sizeof(int)*nthreads);
for (i = 0; i < nthreads; i++)
mp->sync1[i] = 0;
n2 = M * N;
magma_int_t min_mn = min(M, N);
magma_int_t nb = magma_get_zgeqrf_nb(min_mn);
magma_int_t lwork = N*nb;
/* Allocate host memory for the matrix */
TESTING_MALLOC ( h_A , cuDoubleComplex, n2 );
TESTING_MALLOC ( tau , cuDoubleComplex, min_mn);
TESTING_HOSTALLOC( h_R , cuDoubleComplex, n2 );
TESTING_HOSTALLOC(h_work, cuDoubleComplex, lwork );
printf("\n\n");
printf(" M N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("==========================================================\n");
for(i=0; i<10; i++){
if (loop==1){
M = N = min_mn = size[i];
n2 = M*N;
}
flops = FLOPS( (double)M, (double)N ) / 1000000;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &M, h_R, &M );
//magma_zgeqrf(M, N, h_R, M, tau, h_work, lwork, &info);
for(j=0; j<n2; j++)
h_R[j] = h_A[j];
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_qr_init(mp, M, N, h_R, nthreads);
start = get_current_time();
magma_zgeqrf3(context, M, N, h_R, M, tau, h_work, lwork, &info);
end = get_current_time();
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
if (accuracyflag == 1)
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, h_work, &lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of zgeqrf had an illegal value.\n", -info);
cpu_perf = 4.*M*N*min_mn/(3.*1000000*GetTimerValue(start,end));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
double work[1], matnorm = 1.;
cuDoubleComplex mone = MAGMA_Z_NEG_ONE;
magma_int_t one = 1;
if (accuracyflag == 1){
matnorm = lapackf77_zlange("f", &M, &N, h_A, &M, work);
blasf77_zaxpy(&n2, &mone, h_A, &one, h_R, &one);
}
if (accuracyflag == 1){
printf("%5d %5d %6.2f %6.2f %e\n",
M, N, cpu_perf, gpu_perf,
lapackf77_zlange("f", &M, &N, h_R, &M, work) / matnorm);
} else {
printf("%5d %5d %6.2f \n",
M, N, gpu_perf);
}
if (loop != 1)
break;
}
/* Memory clean up */
TESTING_FREE ( h_A );
TESTING_FREE ( tau );
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R );
/* Shut down the MAGMA context */
magma_finalize(context);
}
int main(int argc, char **argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
double flops, magma_perf, cuda_perf, error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
FILE *fp ;
magma_int_t i, lda, Xm, Ym;
magma_int_t M, M0 = 0;
magma_int_t N, N0 = 0;
magma_int_t szeA, szeX, szeY;
magma_int_t istart = 64;
magma_int_t iend = 10240;
magma_int_t incx = 1;
magma_int_t incy = 1;
char trans = MagmaNoTrans;
cuDoubleComplex alpha = MAGMA_Z_MAKE(1., 0.); // MAGMA_Z_MAKE( 1.5, -2.3 );
cuDoubleComplex beta = MAGMA_Z_MAKE(0., 0.); // MAGMA_Z_MAKE( -0.6, 0.8 );
cuDoubleComplex *A, *X, *Y, *Ycublas, *Ymagma;
cuDoubleComplex *dA, *dX, *dY;
if (argc != 1){
for(i=1; i<argc; i++){
if ( strcmp("-n", argv[i]) == 0 ){
N0 = atoi(argv[++i]);
}
else if ( strcmp("-m", argv[i]) == 0 ){
M0 = atoi(argv[++i]);
}
else if (strcmp("-N", argv[i])==0){
trans = MagmaNoTrans;
}
else if (strcmp("-T", argv[i])==0){
trans = MagmaTrans;
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else if (strcmp("-C", argv[i])==0){
trans = MagmaConjTrans;
}
#endif
}
}
if ( (M0 != 0) && (N0 != 0) )
iend = istart + 1;
M = N = iend;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if( trans == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
lda = ((M+31)/32)*32;
szeA = lda*N;
szeX = incx*Xm;
szeY = incy*Ym;
TESTING_MALLOC( A, cuDoubleComplex, szeA );
TESTING_MALLOC( X, cuDoubleComplex, szeX );
TESTING_MALLOC( Y, cuDoubleComplex, szeY );
TESTING_MALLOC( Ycublas, cuDoubleComplex, szeY );
TESTING_MALLOC( Ymagma, cuDoubleComplex, szeY );
TESTING_DEVALLOC( dA, cuDoubleComplex, szeA );
TESTING_DEVALLOC( dX, cuDoubleComplex, szeX );
TESTING_DEVALLOC( dY, cuDoubleComplex, szeY );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &szeA, A );
lapackf77_zlarnv( &ione, ISEED, &szeX, X );
lapackf77_zlarnv( &ione, ISEED, &szeY, Y );
fp = fopen ("results_zgemv.txt", "w") ;
if( fp == NULL ){ printf("Couldn't open output file\n"); exit(1);}
printf("\nUsage: \n");
printf(" testing_zgemv [-N|T|C] [-m %d] [-n %d]\n\n", 1024, 1024);
printf( " m n CUBLAS,Gflop/s MAGMABLAS Gflop/s \"error\"\n"
"==============================================================\n");
fprintf(fp, " m n CUBLAS,Gflop/s MAGMABLAS Gflop/s \"error\"\n"
"==============================================================\n");
for( i=istart; i < iend; i = (int)((i+1)*1.1) )
{
M = N = i;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if( trans == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
lda = ((M+31)/32)*32;
flops = FLOPS( (double)M, (double)N ) / 1000000;
printf( "%5d %5d ", (int) M, (int) N );
fprintf( fp, "%5d %5d ", (int) M, (int) N );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( M, N, A, lda, dA, lda );
magma_zsetvector( Xm, X, incx, dX, incx );
magma_zsetvector( Ym, Y, incy, dY, incy );
/*
* Cublas Version
*/
start = get_current_time();
cublasZgemv( trans, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
end = get_current_time();
magma_zgetvector( Ym, dY, incy, Ycublas, incy );
cuda_perf = flops / GetTimerValue(start, end);
printf( "%11.2f", cuda_perf );
fprintf(fp, "%11.2f", cuda_perf );
/*
* Magma Version
*/
magma_zsetvector( Ym, Y, incy, dY, incy );
start = get_current_time();
magmablas_zgemv( trans, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
end = get_current_time();
magma_zgetvector( Ym, dY, incx, Ymagma, incx );
magma_perf = flops / GetTimerValue(start, end);
printf( "%11.2f", magma_perf );
fprintf(fp, "%11.2f", magma_perf );
/* =====================================================================
Computing the Difference Cublas VS Magma
=================================================================== */
blasf77_zaxpy( &Ym, &c_neg_one, Ymagma, &incy, Ycublas, &incy);
error = lapackf77_zlange( "M", &Ym, &ione, Ycublas, &Ym, work );
#if 0
printf( "\t\t %8.6e", error / (double)Ym );
fprintf( fp, "\t\t %8.6e", error / (double)Ym );
/*
* Blas comparaison
*/
{
char *blastrans = MagmaNoTransStr;
if ( trans == MagmaConjTrans )
blastrans = MagmaConjTransStr;
else if ( trans == MagmaTrans )
blastrans = MagmaTransStr;
blasf77_zcopy( &Ym, Y, &incy, Ycublas, &incy );
blasf77_zgemv( blastrans, &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Ycublas, &incy );
blasf77_zaxpy( &Ym, &c_neg_one, Ymagma, &incy, Ycublas, &incy);
error = lapackf77_zlange( "M", &Ym, &ione, Ycublas, &Ym, work );
}
#endif
printf( "\t\t %8.6e\n", error / (double)Ym );
fprintf( fp, "\t\t %8.6e\n", error / (double)Ym );
}
/* Free Memory */
TESTING_FREE( A );
TESTING_FREE( X );
TESTING_FREE( Y );
TESTING_FREE( Ycublas );
TESTING_FREE( Ymagma );
TESTING_DEVFREE( dA );
TESTING_DEVFREE( dX );
TESTING_DEVFREE( dY );
/* Free device */
TESTING_CUDA_FINALIZE();
return EXIT_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zlarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t M, N, K, size, ldc, ldv, ldt, ldw, nv;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double error, work[1];
// test all combinations of input parameters
const char side[] = { MagmaLeft, MagmaRight };
const char trans[] = { MagmaConjTrans, MagmaNoTrans };
const char direct[] = { MagmaForward, MagmaBackward };
const char storev[] = { MagmaColumnwise, MagmaRowwise };
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
M = opts.msize[i];
N = opts.nsize[i];
K = opts.ksize[i];
if ( M < K || N < K || K <= 0 ) {
printf( "skipping M %d, N %d, K %d; requires M >= K, N >= K, K >= 0.\n", (int) M, (int) N, (int) K );
continue;
}
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
ldc = ((M+31)/32)*32;
ldt = ((K+31)/32)*32;
ldw = (side[iside] == MagmaLeft ? N : M);
// (ldv, nv) get swapped later if rowwise
ldv = (side[iside] == MagmaLeft ? M : N);
nv = K;
// Allocate memory for matrices
magmaDoubleComplex *C, *R, *V, *T, *W;
TESTING_MALLOC( C, magmaDoubleComplex, ldc*N );
TESTING_MALLOC( R, magmaDoubleComplex, ldc*N );
TESTING_MALLOC( V, magmaDoubleComplex, ldv*K );
TESTING_MALLOC( T, magmaDoubleComplex, ldt*K );
TESTING_MALLOC( W, magmaDoubleComplex, ldw*K );
magmaDoubleComplex *dC, *dV, *dT, *dW;
TESTING_DEVALLOC( dC, magmaDoubleComplex, ldc*N );
TESTING_DEVALLOC( dV, magmaDoubleComplex, ldv*K );
TESTING_DEVALLOC( dT, magmaDoubleComplex, ldt*K );
TESTING_DEVALLOC( dW, magmaDoubleComplex, ldw*K );
// C is M x N.
size = ldc*N;
lapackf77_zlarnv( &ione, ISEED, &size, C );
//printf( "C=" ); magma_zprint( M, N, C, ldc );
// V is ldv x nv. See larfb docs for description.
// if column-wise and left, M x K
// if column-wise and right, N x K
// if row-wise and left, K x M
// if row-wise and right, K x N
size = ldv*nv;
lapackf77_zlarnv( &ione, ISEED, &size, V );
if ( storev[istor] == MagmaColumnwise ) {
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_zlaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, &V[(ldv-K)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_zlaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, &V[(nv-K)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_zprint( ldv, nv, V, ldv );
// T is K x K, upper triangular for forward, and lower triangular for backward
magma_int_t k1 = K-1;
size = ldt*K;
lapackf77_zlarnv( &ione, ISEED, &size, T );
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_zlaset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_zprint( K, K, T, ldt );
magma_zsetmatrix( M, N, C, ldc, dC, ldc );
magma_zsetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_zsetmatrix( K, K, T, ldt, dT, ldt );
lapackf77_zlarfb( &side[iside], &trans[itran], &direct[idir], &storev[istor],
&M, &N, &K,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_zprint( M, N, C, ldc );
magma_zlarfb_gpu( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_zgetmatrix( M, N, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_zprint( M, N, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_zlange( "Fro", &M, &N, C, &ldc, work );
size = ldc*N;
blasf77_zaxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_zlange( "Fro", &M, &N, R, &ldc, work ) / error;
printf( "%5d %5d %5d %c %c %c %c %8.2e\n",
(int) M, (int) N, (int) K,
storev[istor], side[iside], direct[idir], trans[itran], error );
TESTING_FREE( C );
TESTING_FREE( R );
TESTING_FREE( V );
TESTING_FREE( T );
TESTING_FREE( W );
TESTING_DEVFREE( dC );
TESTING_DEVFREE( dV );
TESTING_DEVFREE( dT );
TESTING_DEVFREE( dW );
}}}}
printf( "\n" );
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvdx
*/
int main( int argc, char** argv)
{
//#define USE_MGPU
#ifdef USE_MGPU
TESTING_CUDA_INIT_MGPU();
#else
TESTING_CUDA_INIT();
#endif
magma_int_t nrgpu =1;
cuDoubleComplex *h_A, *h_R, *h_B, *h_S, *h_work;
double *rwork, *w1, *w2;
magma_int_t *iwork;
double gpu_time, cpu_time;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t size[4] = {1024,2048,4100,6001};
magma_int_t i, itype, info;
magma_int_t ione = 1, izero = 0;
magma_int_t five = 5;
cuDoubleComplex c_zero = MAGMA_Z_ZERO;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double d_one = 1.;
double d_neg_one = -1.;
double d_ten = 10.;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t il,iu,m1,m2;
double vl,vu;
double fraction_ev = 0;
//const char *uplo = MagmaLowerStr;
char *uplo = (char*)MagmaLowerStr;
//char *uplo = (char*)MagmaUpperStr;
char *jobz = (char*)MagmaVectorsStr;
char range = 'A';
itype = 1;
magma_int_t checkres;
double result[2];
int flagN = 0;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0){
N = atoi(argv[++i]);
if (N>0){
printf(" testing_zhegvdx -N %d\n\n", (int) N);
flagN=1;
}
else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -N %d\n\n", (int) N);
exit(1);
}
}
if (strcmp("-ngpu", argv[i])==0){
nrgpu = atoi(argv[++i]);
if (nrgpu>0){
printf(" testing_zhegvdx -ngpu %d\n\n", (int) nrgpu);
}
else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -ngpu %d\n\n", (int) nrgpu);
exit(1);
}
}
if (strcmp("-itype", argv[i])==0){
itype = atoi(argv[++i]);
if (itype>0 && itype <= 3){
printf(" testing_zhegvdx -itype %d\n\n", (int) itype);
}
else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -itype %d\n\n", (int) itype);
exit(1);
}
}
if (strcmp("-FE", argv[i])==0){
fraction_ev = atof(argv[++i]);
if (fraction_ev > 0 && fraction_ev <= 1){
printf(" testing_zhegvdx -FE %f\n\n", fraction_ev);
}
else {
fraction_ev = 0;
}
}
if (strcmp("-L", argv[i])==0){
uplo = (char*)MagmaLowerStr;
printf(" testing_zhegvdx -L");
}
if (strcmp("-U", argv[i])==0){
uplo = (char*)MagmaUpperStr;
printf(" testing_zhegvdx -U");
}
}
} else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -L/U -N %d -itype %d\n\n", 1024, 1);
}
if(!flagN)
N = size[3];
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
n2 = N * N;
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, cuDoubleComplex, n2);
TESTING_MALLOC( h_B, cuDoubleComplex, n2);
TESTING_MALLOC( w1, double , N);
TESTING_MALLOC( w2, double , N);
TESTING_HOSTALLOC(h_R, cuDoubleComplex, n2);
TESTING_HOSTALLOC(h_S, cuDoubleComplex, n2);
magma_int_t nb = magma_get_zhetrd_nb(N);
magma_int_t lwork = magma_zbulge_get_lq2(N) + 2*N + N*N;
magma_int_t lrwork = 1 + 5*N +2*N*N;
magma_int_t liwork = 3 + 5*N;
TESTING_HOSTALLOC(h_work, cuDoubleComplex, lwork);
TESTING_HOSTALLOC( rwork, double, lrwork);
TESTING_MALLOC( iwork, magma_int_t, liwork);
printf(" N M GPU Time(s) \n");
printf("==========================\n");
for(i=0; i<4; i++){
if (!flagN){
N = size[i];
n2 = N*N;
}
if (fraction_ev == 0){
il = N / 10;
iu = N / 5+il;
}
else {
il = 1;
iu = (int)(fraction_ev*N);
if (iu < 1) iu = 1;
}
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
//lapackf77_zlatms( &N, &N, "U", ISEED, "P", w1, &five, &d_ten,
// &d_one, &N, &N, uplo, h_B, &N, h_work, &info);
//lapackf77_zlaset( "A", &N, &N, &c_zero, &c_one, h_B, &N);
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_Z_SET2REAL( h_B[i*N+i], ( MAGMA_Z_REAL(h_B[i*N+i]) + 1.*N ) );
}
}
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
#ifdef USE_MGPU
magma_zhegvdx_2stage_m(nrgpu, itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#else
magma_zhegvdx_2stage(itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#endif
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
#ifdef USE_MGPU
magma_zhegvdx_2stage_m(nrgpu, itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#else
magma_zhegvdx_2stage(itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#endif
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
/* =====================================================================
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 |
=================================================================== */
double temp1, temp2;
cuDoubleComplex *tau;
result[0] = 1.;
result[0] /= lapackf77_zlanhe("1",uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_zlange("1",&N , &m1, h_R, &N, rwork);
if (itype == 1){
blasf77_zhemm("L", 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", 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 (itype == 2){
blasf77_zhemm("L", 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", 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 (itype == 3){
blasf77_zhemm("L", 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", 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(itype, 'N', range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&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 / temp1;
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %6.2f\n",
(int) N, (int) m1, gpu_time);
if ( checkres ){
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
printf("(2) | D(w/ Z)-D(w/o Z)|/ |D| = %e\n\n", result[1]);
}
if (flagN)
break;
}
cudaSetDevice(0);
/* Memory clean up */
TESTING_FREE( h_A);
TESTING_FREE( h_B);
TESTING_FREE( w1);
TESTING_FREE( w2);
TESTING_HOSTFREE( rwork);
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
TESTING_HOSTFREE( h_S);
/* Shutdown */
#ifdef USE_MGPU
TESTING_CUDA_FINALIZE_MGPU();
#else
TESTING_CUDA_FINALIZE();
#endif
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrs
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double gpu_error, cpu_error, matnorm, work[1];
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_A2, *h_B, *h_X, *h_R, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_A, *d_B;
magma_int_t M, N, n2, nrhs, lda, ldb, ldda, lddb, min_mn, max_mn, nb, info;
magma_int_t lworkgpu, lhwork, lhwork2;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol = opts.tolerance * lapackf77_dlamch("E");
nrhs = opts.nrhs;
printf(" ||b-Ax|| / (N||A||)\n");
printf(" M N NRHS CPU GFlop/s (sec) GPU GFlop/s (sec) CPU GPU \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];
if ( M < N ) {
printf( "skipping M=%d, N=%d because M < N is not yet supported.\n", (int) M, (int) N );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
n2 = lda*N;
ldda = ((M+31)/32)*32;
lddb = ((max_mn+31)/32)*32;
nb = magma_get_zgeqrf_nb(M);
gflops = (FLOPS_ZGEQRF( M, N ) + FLOPS_ZGEQRS( M, N, nrhs )) / 1e9;
// query for workspace size
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
lhwork = -1;
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
lhwork2 = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
lhwork = -1;
lapackf77_zunmqr( MagmaLeftStr, MagmaConjTransStr,
&M, &nrhs, &min_mn, h_A, &lda, tau,
h_X, &ldb, tmp, &lhwork, &info);
lhwork = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
lhwork = max( max( lhwork, lhwork2 ), lworkgpu );
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC( h_X, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC( h_R, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC( h_work, magmaDoubleComplex, lhwork );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*nrhs );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
n2 = M*nrhs;
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
// make consistent RHS
//n2 = N*nrhs;
//lapackf77_zlarnv( &ione, ISEED, &n2, h_X );
//blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
// &c_one, h_A, &lda,
// h_X, &ldb,
// &c_zero, h_B, &ldb );
//lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_zgels3_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_work, lworkgpu, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get the solution in h_X
magma_zgetmatrix( N, nrhs, d_B, lddb, h_X, ldb );
// compute the residual
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
matnorm = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_zgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_work, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb);
cpu_error = lapackf77_zlange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*matnorm);
gpu_error = lapackf77_zlange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*matnorm);
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, (int) nrhs,
cpu_perf, cpu_time, gpu_perf, gpu_time, cpu_error, gpu_error );
printf("%s\n", (gpu_error < tol ? "" : " failed"));
status |= ! (gpu_error < tol);
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_A2 );
TESTING_FREE( h_B );
TESTING_FREE( h_X );
TESTING_FREE( h_R );
TESTING_FREE( h_work );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
magmaDoubleComplex *hA, *hR;
magmaDoubleComplex_ptr dA;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] =
{ 1024, 2048, 3072, 4032, 5184, 6048, 7200, 8064, 8928, 10560 };
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_zpotrf_gpu -N %d\n\n", 1024);
}
/* Initialize */
magma_queue_t queue;
magma_device_t device;
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( &device, 1, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device, &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
/* Allocate memory for the largest matrix */
N = size[9];
n2 = N * N;
ldda = ((N+31)/32) * 32;
TESTING_MALLOC( hA, magmaDoubleComplex, n2 );
TESTING_MALLOC_HOST( 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 ) {
MAGMA_Z_SET2REAL( hA(i,i), MAGMA_Z_REAL(hA(i,i)) + N );
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, 0, lda, dA, 0, ldda, queue );
magma_zpotrf_gpu( MagmaUpper, N, dA, 0, ldda, &info, queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( N, N, hA, 0, lda, dA, 0, ldda, queue );
gpu_time = get_time();
magma_zpotrf_gpu( MagmaUpper, N, dA, 0, ldda, &info, queue );
gpu_time = get_time() - gpu_time;
if (info != 0)
printf( "magma_zpotrf had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = get_time();
lapackf77_zpotrf( MagmaUpperStr, &N, hA, &lda, &info );
cpu_time = get_time() - 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, 0, lda, queue );
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( hA );
TESTING_FREE_HOST( hR );
TESTING_FREE_DEV( dA );
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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;
magma_int_t N, n2, lda, 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;
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("ngpu %d, uplo %c\n", (int) opts.ngpu, 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 i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
gflops = FLOPS_ZPOTRF( N ) / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
/* 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
=================================================================== */
gpu_time = magma_wtime();
magma_zpotrf( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( &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
=================================================================== */
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 ? "" : " failed") );
status |= ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
real_Double_t mgpu_time;
magmaDoubleComplex *h_A, *h_Ainit, *h_B, *h_Binit, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
magma_int_t lrwork;
#endif
double *w1, result;
magma_int_t *iwork;
magma_int_t N, n2, info, lwork, liwork;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_timestr_t start, end;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
char jobz = opts.jobz;
int checkres = opts.check;
char range = 'A';
char uplo = opts.uplo;
magma_int_t itype = opts.itype;
double f = opts.fraction;
if (f != 1)
range='I';
if ( checkres && jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
jobz = MagmaVec;
}
printf("using: nrgpu = %d, itype = %d, jobz = %c, range = %c, uplo = %c, checkres = %d, fraction = %6.4f\n",
(int) opts.ngpu, (int) itype, jobz, range, uplo, (int) checkres, f);
printf(" N M nr GPU MGPU Time(s) \n");
printf("====================================\n");
magma_int_t threads = magma_get_numthreads();
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = magma_zbulge_get_lq2(N, threads) + 2*N + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = magma_zbulge_get_lq2(N, threads) + 1 + 6*N + 2*N*N;
#endif
liwork = 3 + 5*N;
//magma_int_t NB = 96;//magma_bulge_get_nb(N);
//magma_int_t sizvblg = magma_zbulge_get_lq2(N, threads);
//magma_int_t siz = max(sizvblg,n2)+2*(N*NB+N)+24*N;
/* Allocate host memory for the matrix */
TESTING_HOSTALLOC( h_A, magmaDoubleComplex, n2);
TESTING_HOSTALLOC( h_B, magmaDoubleComplex, n2);
TESTING_MALLOC( w1, double , N);
TESTING_HOSTALLOC(h_work, magmaDoubleComplex, lwork);
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTALLOC( rwork, double, lrwork);
#endif
TESTING_MALLOC( iwork, magma_int_t, liwork);
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
{
for(int i=0; i<N; i++) {
MAGMA_Z_SET2REAL( h_B[i*N+i], ( MAGMA_Z_REAL(h_B[i*N+i]) + 1.*N ) );
MAGMA_Z_SET2REAL( h_A[i*N+i], MAGMA_Z_REAL(h_A[i*N+i]) );
}
}
if((opts.warmup)||( checkres )){
TESTING_MALLOC(h_Ainit, magmaDoubleComplex, n2);
TESTING_MALLOC(h_Binit, magmaDoubleComplex, n2);
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_Binit, &N );
}
magma_int_t m1 = 0;
double vl = 0;
double vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (range == 'I'){
il = 1;
iu = (int) (f*N);
}
if(opts.warmup){
// ==================================================================
// Warmup using MAGMA. I prefer to use smalltest to warmup A-
// ==================================================================
magma_zhegvdx_2stage_m(opts.ngpu, itype, jobz, range, uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
start = get_current_time();
magma_zhegvdx_2stage_m(opts.ngpu, itype, jobz, range, uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
end = get_current_time();
mgpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
// ===================================================================
// Check the results following the LAPACK's [zc]hegvdx routine.
// A x = lambda B x is solved
// and the following 3 tests computed:
// (1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
// | A B Z - Z D | / ( |A||Z| N ) (itype = 2)
// | B A Z - Z D | / ( |A||Z| N ) (itype = 3)
// ===================================================================
#if defined(PRECISION_d) || defined(PRECISION_s)
double *rwork = h_work + N*N;
#endif
result = 1.;
result /= lapackf77_zlanhe("1",&uplo, &N, h_Ainit, &N, rwork);
result /= lapackf77_zlange("1",&N , &m1, h_A, &N, rwork);
if (itype == 1){
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_zhemm("L", &uplo, &N, &m1, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result *= lapackf77_zlange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (itype == 2){
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Binit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_zlange("1", &N, &m1, h_A, &N, rwork)/N;
}
else if (itype == 3){
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_zlange("1", &N, &m1, h_A, &N, rwork)/N;
}
}
// ===================================================================
// Print execution time
// ===================================================================
printf("%5d %5d %2d %6.2f\n",
(int) N, (int) m1, (int) opts.ngpu, mgpu_time);
if ( checkres ){
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
}
TESTING_HOSTFREE( h_A);
TESTING_HOSTFREE( h_B);
TESTING_FREE( w1);
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTFREE( rwork);
#endif
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
if((opts.warmup)||( checkres )){
TESTING_FREE( h_Ainit);
TESTING_FREE( h_Binit);
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE_MGPU();
return 0;
}