/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqlf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const double d_neg_one = MAGMA_D_NEG_ONE;
const double d_one = MAGMA_D_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double Anorm, error=0, error2=0;
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 ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |L - Q^H*A| |I - Q^H*Q|\n");
printf("===============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
nb = magma_get_zgeqlf_nb(M);
gflops = FLOPS_ZGEQLF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgeqlf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, N*nb );
lwork = max( lwork, 2*nb*nb);
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
/* 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 ));
/* =====================================================================
Check the result, following zqlt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
=================================================================== */
if ( opts.check ) {
magma_int_t ldq = M;
magma_int_t ldl = min_mn;
magmaDoubleComplex *Q, *L;
double *work;
TESTING_MALLOC_CPU( Q, magmaDoubleComplex, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( L, magmaDoubleComplex, ldl*N ); // K by N
TESTING_MALLOC_CPU( work, double, min_mn );
// copy M by K matrix V to Q (copying diagonal, which isn't needed) and
// copy K by N matrix L
lapackf77_zlaset( "Full", &min_mn, &N, &c_zero, &c_zero, L, &ldl );
if ( M >= N ) {
// for M=5, N=3: A = [ V V V ] <= V full block (M-N by K)
// K=N [ V V V ]
// [ ----- ]
// [ L V V ] <= V triangle (N by K, copying diagonal too)
// [ L L V ] <= L triangle (K by N)
// [ L L L ]
magma_int_t M_N = M - N;
lapackf77_zlacpy( "Full", &M_N, &min_mn, h_R, &lda, Q, &ldq );
lapackf77_zlacpy( "Upper", &N, &min_mn, &h_R[M_N], &lda, &Q[M_N], &ldq );
lapackf77_zlacpy( "Lower", &min_mn, &N, &h_R[M_N], &lda, L, &ldl );
}
else {
// for M=3, N=5: A = [ L L | L V V ] <= V triangle (K by K)
// K=M [ L L | L L V ] <= L triangle (K by M)
// [ L L | L L L ]
// ^^^============= L full block (K by N-M)
magma_int_t N_M = N - M;
lapackf77_zlacpy( "Upper", &M, &min_mn, &h_R[N_M*lda], &lda, Q, &ldq );
lapackf77_zlacpy( "Full", &min_mn, &N_M, h_R, &lda, L, &ldl );
lapackf77_zlacpy( "Lower", &min_mn, &M, &h_R[N_M*lda], &lda, &L[N_M*ldl], &ldl );
}
// generate M by K matrix Q, where K = min(M,N)
lapackf77_zungql( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// error = || L - Q^H*A || / (N * ||A||)
blasf77_zgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, L, &ldl );
Anorm = lapackf77_zlange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_zlange( "1", &min_mn, &N, L, &ldl, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set L = I (K by K identity), then L = I - Q^H*Q
// error = || I - Q^H*Q || / N
lapackf77_zlaset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, L, &ldl );
blasf77_zherk( "Upper", "Conj", &min_mn, &M, &d_neg_one, Q, &ldq, &d_one, L, &ldl );
error2 = lapackf77_zlanhe( "1", "Upper", &min_mn, L, &ldl, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( L ); L = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.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 ));
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check ) {
bool okay = (error < tol && error2 < tol);
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error;
magmaFloatComplex *h_A, *h_R;
magmaFloatComplex *d_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||PA-LU||/(||A||*N)\n");
printf("=========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGETRF( M, N ) / 1e9;
if ( N > 512 ) {
printf( "%5d %5d skipping because cgetf2 does not support N > 512\n", (int) M, (int) N );
continue;
}
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* 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 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgetrf(&M, &N, h_A, &lda, ipiv, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgetf2_gpu( M, N, d_A, ldda, ipiv, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgetf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000. );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time*1000. );
}
if ( opts.check ) {
magma_cgetmatrix( M, N, d_A, ldda, h_A, lda );
error = get_LU_error( M, N, h_R, lda, h_A, ipiv );
printf(" %8.2e %s\n", error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const float d_neg_one = MAGMA_D_NEG_ONE;
const float d_one = MAGMA_D_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
const float c_one = MAGMA_S_ONE;
const float c_zero = MAGMA_S_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float Anorm, error=0, error2=0;
float *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloat_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
opts.parse_opts( argc, argv );
magma_int_t status = 0;
float tol = opts.tolerance * lapackf77_slamch("E");
// version 3 can do either check
if (opts.check == 1 && opts.version == 1) {
opts.check = 2;
printf( "%% version 1 requires check 2 (solve A*x=b)\n" );
}
if (opts.check == 2 && opts.version == 2) {
opts.check = 1;
printf( "%% version 2 requires check 1 (R - Q^H*A)\n" );
}
printf( "%% version %d\n", (int) opts.version );
if ( opts.check == 1 ) {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |R - Q^H*A| |I - Q^H*Q|\n");
printf("%%==============================================================================\n");
}
else {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |b - A*x|\n");
printf("%%===============================================================\n");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min( M, N );
lda = M;
n2 = lda*N;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
nb = magma_get_sgeqrf_nb( M, N );
gflops = FLOPS_SGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgeqrf( &M, &N, NULL, &M, NULL, tmp, &lwork, &info );
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_work, float, lwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
if ( opts.version == 1 || opts.version == 3 ) {
size = (2*min(M, N) + magma_roundup( N, 32 ) )*nb;
TESTING_MALLOC_DEV( dT, float, size );
magmablas_slaset( MagmaFull, size, 1, c_zero, c_zero, dT, size );
}
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
nb = magma_get_sgeqrf_nb( M, N );
gpu_time = magma_wtime();
if ( opts.version == 1 ) {
// stores dT, V blocks have zeros, R blocks inverted & stored in dT
magma_sgeqrf_gpu( M, N, d_A, ldda, tau, dT, &info );
}
else if ( opts.version == 2 ) {
// LAPACK complaint arguments
magma_sgeqrf2_gpu( M, N, d_A, ldda, tau, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// stores dT, V blocks have zeros, R blocks stored in dT
magma_sgeqrf3_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
return -1;
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( opts.check == 1 && (opts.version == 2 || opts.version == 3) ) {
if ( opts.version == 3 ) {
// copy diagonal blocks of R back to A
for( int i=0; i < min_mn-nb; i += nb ) {
magma_int_t ib = min( min_mn-i, nb );
magmablas_slacpy( MagmaUpper, ib, ib, &dT[min_mn*nb + i*nb], nb, &d_A[ i + i*ldda ], ldda );
}
}
/* =====================================================================
Check the result, following zqrt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
Only for version 2, which has LAPACK complaint output.
Or for version 3, after restoring diagonal blocks of A above.
=================================================================== */
magma_sgetmatrix( M, N, d_A, ldda, h_R, lda );
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
float *Q, *R;
float *work;
TESTING_MALLOC_CPU( Q, float, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, float, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, float, min_mn );
// generate M by K matrix Q, where K = min(M,N)
lapackf77_slacpy( "Lower", &M, &min_mn, h_R, &lda, Q, &ldq );
lapackf77_sorgqr( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy K by N matrix R
lapackf77_slaset( "Lower", &min_mn, &N, &c_zero, &c_zero, R, &ldr );
lapackf77_slacpy( "Upper", &min_mn, &N, h_R, &lda, R, &ldr );
// error = || R - Q^H*A || / (N * ||A||)
blasf77_sgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, R, &ldr );
Anorm = lapackf77_slange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_slange( "1", &min_mn, &N, R, &ldr, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set R = I (K by K identity), then R = I - Q^H*Q
// error = || I - Q^H*Q || / N
lapackf77_slaset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, R, &ldr );
blasf77_ssyrk( "Upper", "Conj", &min_mn, &M, &d_neg_one, Q, &ldq, &d_one, R, &ldr );
error2 = safe_lapackf77_slansy( "1", "Upper", &min_mn, R, &ldr, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
else if ( opts.check == 2 && M >= N && (opts.version == 1 || opts.version == 3) ) {
/* =====================================================================
Check the result by solving consistent linear system, A*x = b.
Only for versions 1 & 3 with M >= N.
=================================================================== */
magma_int_t lwork2;
float *x, *b, *hwork;
magmaFloat_ptr d_B;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, float, N );
TESTING_MALLOC_CPU( b, float, M );
lapackf77_slarnv( &ione, ISEED, &N, x );
blasf77_sgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, float, M );
magma_ssetvector( M, b, 1, d_B, 1 );
if ( opts.version == 1 ) {
// allocate hwork
magma_sgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork2 = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lwork2 );
// solve linear system
magma_sgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork2, &info );
if (info != 0) {
printf("magma_sgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_sgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork2 = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lwork2 );
// solve linear system
magma_sgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork2, &info );
if (info != 0) {
printf("magma_sgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
return -1;
}
magma_sgetvector( N, d_B, 1, x, 1 );
// compute r = Ax - b, saved in b
blasf77_sgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (max(m,n)*|A|*|x|)
float norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_slange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_slange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_slange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (max(M,N) * norm_A * norm_x);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_sgeqrf( &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_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check == 1 ) {
bool okay = (error < tol && error2 < tol);
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else if ( opts.check == 2 ) {
if ( M >= N ) {
bool okay = (error < tol);
status += ! okay;
printf( "%10.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf( "(error check only for M >= N)\n" );
}
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version == 1 || opts.version == 3 ) {
TESTING_FREE_DEV( dT );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dsyevd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
double *h_A, *h_R, *h_work;
double *w1, *w2;
magma_int_t *iwork;
magma_int_t N, n2, info, lwork, liwork, lda, aux_iwork[1];
magma_int_t izero = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double result[3], eps, aux_work[1];
eps = lapackf77_dlamch( "E" );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: jobz = %s, uplo = %s\n",
lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo));
printf(" N CPU Time (sec) GPU Time (sec)\n");
printf("=======================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
lda = N;
// query for workspace sizes
magma_dsyevd( opts.jobz, opts.uplo,
N, NULL, lda, NULL,
aux_work, -1,
aux_iwork, -1,
opts.queue, &info );
lwork = (magma_int_t) aux_work[0];
liwork = aux_iwork[0];
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, double, N*lda );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, double, N*lda );
TESTING_MALLOC_PIN( h_work, double, lwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* warm up run */
if ( opts.warmup ) {
magma_dsyevd( opts.jobz, opts.uplo,
N, h_R, lda, w1,
h_work, lwork,
iwork, liwork,
opts.queue, &info );
if (info != 0)
printf("magma_dsyevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dsyevd( opts.jobz, opts.uplo,
N, h_R, lda, w1,
h_work, lwork,
iwork, liwork,
opts.queue, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_dsyevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
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 |
=================================================================== */
double temp1, temp2;
// tau=NULL is unused since itype=1
lapackf77_dsyt21( &ione, lapack_uplo_const(opts.uplo), &N, &izero,
h_A, &lda,
w1, h_work,
h_R, &lda,
h_R, &lda,
NULL, h_work, &result[0] );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dsyevd( MagmaNoVec, opts.uplo,
N, h_R, lda, w2,
h_work, lwork,
iwork, liwork,
opts.queue, &info );
if (info != 0)
printf("magma_dsyevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for( int j=0; j<N; j++ ) {
temp1 = max(temp1, fabs(w1[j]));
temp1 = max(temp1, fabs(w2[j]));
temp2 = max(temp2, fabs(w1[j]-w2[j]));
}
result[2] = temp2 / (((double)N)*temp1);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dsyevd( lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
&N, h_A, &lda, w2,
h_work, &lwork,
iwork, &liwork,
&info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_dsyevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf("%5d %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time);
}
else {
printf("%5d --- %7.2f\n",
(int) N, gpu_time);
}
/* =====================================================================
Print execution time
=================================================================== */
if ( opts.check ) {
printf("Testing the factorization A = U S U' for correctness:\n");
printf("(1) | A - U S U' | / (|A| N) = %8.2e %s\n", result[0]*eps, (result[0]*eps < tol ? "ok" : "failed") );
printf("(2) | I - U'U | / N = %8.2e %s\n", result[1]*eps, (result[1]*eps < tol ? "ok" : "failed") );
printf("(3) | S(w/ U) - S(w/o U) | / |S| = %8.2e %s\n\n", result[2] , (result[2] < tolulp ? "ok" : "failed") );
status += ! (result[0]*eps < tol && result[1]*eps < tol && result[2] < tolulp);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgehrd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A, *h_R, *h_Q, *h_work, *tau, *twork;
magmaFloatComplex_ptr dT;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
#endif
float eps, result[2];
magma_int_t N, n2, lda, nb, lwork, ltwork, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
eps = lapackf77_slamch( "E" );
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) |A-QHQ'|/N|A| |I-QQ'|/N\n");
printf("=========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
nb = magma_get_cgehrd_nb(N);
/* We suppose the magma nb is bigger than lapack nb */
lwork = N*nb;
gflops = FLOPS_CGEHRD( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, N );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, nb*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgehrd( N, ione, N, h_R, lda, tau, h_work, lwork, dT, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgehrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.check ) {
ltwork = 2*(N*N);
TESTING_MALLOC_PIN( h_Q, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( twork, magmaFloatComplex, ltwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_CPU( rwork, float, N );
#endif
lapackf77_clacpy(MagmaUpperLowerStr, &N, &N, h_R, &lda, h_Q, &lda);
for( int j = 0; j < N-1; ++j )
for( int i = j+2; i < N; ++i )
h_R[i+j*lda] = MAGMA_C_ZERO;
magma_cunghr(N, ione, N, h_Q, lda, tau, dT, nb, &info);
if (info != 0) {
printf("magma_cunghr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
exit(1);
}
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_chst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, rwork, result);
#else
lapackf77_chst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, result);
#endif
TESTING_FREE_PIN( h_Q );
TESTING_FREE_CPU( twork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_CPU( rwork );
#endif
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgehrd(&N, &ione, &N, h_R, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgehrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Print performance and error.
=================================================================== */
if ( opts.lapack ) {
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f)",
(int) N, gpu_perf, gpu_time );
}
if ( opts.check ) {
printf(" %8.2e %8.2e %s\n",
result[0]*eps, result[1]*eps,
( ( (result[0]*eps < tol) && (result[1]*eps < tol) ) ? "ok" : "failed") );
status += ! (result[0]*eps < tol);
status += ! (result[1]*eps < tol);
}
else {
printf(" --- ---\n");
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( tau );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = 2. * opts.tolerance * lapackf77_dlamch("E");
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
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, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, N*nb );
lwork = max( lwork, 2*nb*nb);
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
/* 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 %s\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
float *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_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_SGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_work, float, lwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( dtau, float, min_mn );
TESTING_MALLOC_DEV( dwork, float, min_mn );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
if ( opts.warmup ) {
magma_sgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_sgeqr2_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_sgeqr2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sgeqrf(&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_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_sgetmatrix( M, N, d_A, ldda, h_R, M );
error = lapackf77_slange("f", &M, &N, h_A, &lda, work);
blasf77_saxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegst
*/
int main( int argc, char** argv)
{
TESTING_INIT();
// Constants
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
// Local variables
real_Double_t gpu_time, cpu_time;
magmaFloatComplex *h_A, *h_B, *h_R;
magmaFloatComplex_ptr d_A, d_B;
float Anorm, error, work[1];
magma_int_t N, n2, lda, ldda, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% itype N CPU time (sec) GPU time (sec) |R| \n");
printf("%%=======================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = magma_roundup( lda, opts.align );
n2 = N*lda;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, lda*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, ldda*N );
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clarnv( &ione, ISEED, &n2, h_B );
magma_cmake_hermitian( N, h_A, lda );
magma_cmake_hpd( N, h_B, lda );
magma_cpotrf( opts.uplo, N, h_B, lda, &info );
if (info != 0) {
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
magma_csetmatrix( N, N, h_B, lda, d_B, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_chegst_gpu( opts.itype, opts.uplo, N, d_A, ldda, d_B, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_chegst_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_chegst( &opts.itype, lapack_uplo_const(opts.uplo),
&N, h_A, &lda, h_B, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0) {
printf("lapackf77_chegst returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
Anorm = safe_lapackf77_clanhe("f", lapack_uplo_const(opts.uplo), &N, h_A, &lda, work );
error = safe_lapackf77_clanhe("f", lapack_uplo_const(opts.uplo), &N, h_R, &lda, work )
/ Anorm;
bool okay = (error < tol);
status += ! okay;
printf("%3d %5d %7.2f %7.2f %8.2e %s\n",
(int) opts.itype, (int) N, cpu_time, gpu_time,
error, (okay ? "ok" : "failed"));
}
else {
printf("%3d %5d --- %7.2f\n",
(int) opts.itype, (int) N, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
float *h_R = NULL, *h_P = NULL;
magmaFloat_ptr d_lA[MagmaMaxSubs * MagmaMaxGPUs];
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] =
{ 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 9000, 10000 };
magma_int_t i, j, k, check = 0, info;
float mz_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t num_gpus0 = 1, num_gpus, num_subs0 = 1, num_subs, tot_subs, flag = 0;
magma_int_t nb, n_local, nk;
magma_uplo_t uplo = MagmaLower;
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;
flag = 1;
}
}
if(strcmp("-NGPU", argv[i]) == 0)
num_gpus0 = atoi(argv[++i]);
if(strcmp("-NSUB", argv[i]) == 0)
num_subs0 = atoi(argv[++i]);
if(strcmp("-UPLO", argv[i]) == 0)
uplo = (strcmp("L", argv[++i]) == 0 ? MagmaLower : MagmaUpper);
if(strcmp("-check", argv[i]) == 0)
check = 1;
}
}
/* Initialize */
magma_queue_t queues[2*MagmaMaxGPUs];
magma_device_t devices[ MagmaMaxGPUs ];
magma_int_t num = 0;
magma_int_t err;
magma_init();
err = magma_getdevices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_getdevices failed: %d\n", (int) err );
exit(-1);
}
for(i=0;i<num_gpus0;i++){
err = magma_queue_create( devices[i], &queues[2*i] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( devices[i], &queues[2*i+1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
}
printf("\nUsing %d GPUs:\n", num_gpus0);
printf(" testing_spotrf_msub -N %d -NGPU %d -NSUB %d -UPLO %c %s\n\n", size[0], num_gpus0,num_subs0,
(uplo == MagmaLower ? 'L' : 'U'),(check == 1 ? "-check" : " "));
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;
gflops = FLOPS_SPOTRF( N ) / 1e9;;
nb = magma_get_spotrf_nb(N);
if (num_subs0*num_gpus0 > N/nb) {
num_gpus = N/nb;
num_subs = 1;
if(N%nb != 0) num_gpus ++;
printf("too many GPUs for the matrix size, using %d GPUs\n", (int)num_gpus);
} else {
num_gpus = num_gpus0;
num_subs = num_subs0;
}
tot_subs = num_subs * num_gpus;
/* Allocate host memory for the matrix */
#ifdef USE_PINNED_CLMEMORY
cl_mem buffer1 = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, n2*sizeof(float), NULL, NULL);
cl_mem buffer2 = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, lda*nb*sizeof(float), NULL, NULL);
for (k=0; k<num_gpus; k++) {
h_R = (float*)clEnqueueMapBuffer(queues[2*k], buffer1, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
n2*sizeof(float), 0, NULL, NULL, NULL);
h_P = (float*)clEnqueueMapBuffer(queues[2*k], buffer2, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
lda*nb*sizeof(float), 0, NULL, NULL, NULL);
}
#else
TESTING_MALLOC_PIN( h_P, float, lda*nb );
TESTING_MALLOC_PIN( h_R, float, n2 );
#endif
/* Initialize the matrix */
init_matrix( N, h_R, lda );
/* Allocate GPU memory */
if (uplo == MagmaUpper) {
ldda = ((N+nb-1)/nb)*nb;
n_local = ((N+nb*tot_subs-1)/(nb*tot_subs))*nb;
} else {
ldda = ((N+nb*tot_subs-1)/(nb*tot_subs))*nb;
n_local = ((N+nb-1)/nb)*nb;
}
for (j=0; j<tot_subs; j++) {
TESTING_MALLOC_DEV( d_lA[j], float, n_local*ldda );
}
/* Warm up to measure the performance */
/* distribute matrix to gpus */
if (uplo == MagmaUpper) {
for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_ssetmatrix( j+nk, nk,
&h_R[j*lda], lda,
d_lA[k], j/(nb*tot_subs)*nb*ldda, ldda,
queues[2*(k%num_gpus)]);
}
} else {
for (j=0; j<N; j+=nb) {
nk = min(nb, N-j);
for (magma_int_t kk = 0; kk<tot_subs; kk++) {
magma_int_t mk = 0;
for (magma_int_t ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
magma_int_t mii = min(nb, N-ii);
lapackf77_slacpy( MagmaFullStr, &mii, &nk, &h_R[ii+j*lda], &lda, &h_P[mk], &lda );
mk += mii;
}
k = ((j+kk*nb)/nb)%tot_subs;
if (mk > 0 && nk > 0) {
magma_ssetmatrix( mk, nk,
h_P, lda,
d_lA[k], j*ldda+(j+kk*nb)/(nb*tot_subs)*nb, ldda,
queues[2*(k%num_gpus)]);
}
}
}
/*for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_ssetmatrix( nk, j+nk, &h_R[j], lda,
d_lA[k], j/(nb*tot_subs)*nb, ldda,
queues[2*(k%num_gpus)]);
}*/
}
magma_spotrf_msub( num_subs, num_gpus, uplo, N, d_lA, 0, ldda, queues, &info );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* distribute matrix to gpus */
if (uplo == MagmaUpper) {
for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_ssetmatrix( j+nk, nk,
&h_R[j*lda], lda,
d_lA[k], j/(nb*tot_subs)*nb*ldda, ldda,
queues[2*(k%num_gpus)]);
}
} else {
for (j=0; j<N; j+=nb) {
nk = min(nb, N-j);
for (magma_int_t kk = 0; kk<tot_subs; kk++) {
magma_int_t mk = 0;
for (magma_int_t ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
magma_int_t mii = min(nb, N-ii);
lapackf77_slacpy( MagmaFullStr, &mii, &nk, &h_R[ii+j*lda], &lda, &h_P[mk], &lda );
mk += mii;
}
k = ((j+kk*nb)/nb)%tot_subs;
if (mk > 0 && nk > 0) {
magma_ssetmatrix( mk, nk,
h_P, lda,
d_lA[k], j*ldda+(j+kk*nb)/(nb*tot_subs)*nb, ldda,
queues[2*(k%num_gpus)]);
}
}
}
/*for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_ssetmatrix( nk, j+nk, &h_R[j], lda,
d_lA[k], (j/(nb*tot_subs)*nb), ldda,
queues[2*(k%num_gpus)]);
}*/
}
gpu_time = magma_wtime();
magma_spotrf_msub( num_subs, num_gpus, uplo, N, d_lA, 0, ldda, queues, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf( "magma_spotrf had error %d.\n", info );
/* gather matrix from gpus */
if (uplo==MagmaUpper) {
for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_sgetmatrix( j+nk, nk,
d_lA[k], j/(nb*tot_subs)*nb*ldda, ldda,
&h_R[j*lda], lda, queues[2*(k%num_gpus)]);
}
} else {
for (j=0; j<N; j+=nb) {
nk = min(nb, N-j);
for (magma_int_t kk = 0; kk<tot_subs; kk++) {
k = ((j+kk*nb)/nb)%tot_subs;
magma_int_t mk = 0;
mk = 0;
for (magma_int_t ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
mk += min(nb, N-ii);
}
if (mk > 0 && nk > 0) {
magma_sgetmatrix( mk, nk,
d_lA[k], j*ldda+(j+kk*nb)/(nb*tot_subs)*nb, ldda,
h_P, lda,
queues[2*(k%num_gpus)]);
}
mk = 0;
for (magma_int_t ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
magma_int_t mii = min(nb, N-ii);
lapackf77_slacpy( MagmaFullStr, &mii, &nk, &h_P[mk], &lda, &h_R[ii+j*lda], &lda );
mk += mii;
}
}
}
/*for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_sgetmatrix( nk, j+nk,
d_lA[k], (j/(nb*tot_subs)*nb), ldda,
&h_R[j], lda, queues[2*(k%num_gpus)] );
}*/
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if (check == 1) {
float work[1], matnorm, diffnorm;
float *h_A;
TESTING_MALLOC_PIN( h_A, float, n2 );
init_matrix( N, h_A, lda );
cpu_time = magma_wtime();
if (uplo == MagmaLower) {
lapackf77_spotrf( MagmaLowerStr, &N, h_A, &lda, &info );
} else {
lapackf77_spotrf( MagmaUpperStr, &N, h_A, &lda, &info );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf( "lapackf77_spotrf had error %d.\n", info );
/* =====================================================================
Check the result compared to LAPACK
|R_magma - R_lapack| / |R_lapack|
=================================================================== */
matnorm = lapackf77_slange("f", &N, &N, h_A, &lda, work);
blasf77_saxpy(&n2, &mz_one, h_A, &ione, h_R, &ione);
diffnorm = lapackf77_slange("f", &N, &N, h_R, &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 );
TESTING_FREE_PIN( h_A );
} else {
printf( "%5d - - (- -) %6.2f (%6.2f) - -\n",
N, gpu_perf, gpu_time );
}
// free memory
#ifdef USE_PINNED_CLMEMORY
for (k=0; k<num_gpus; k++) {
clEnqueueUnmapMemObject(queues[2*k], buffer1, h_R, 0, NULL, NULL);
clEnqueueUnmapMemObject(queues[2*k], buffer2, h_P, 0, NULL, NULL);
}
clReleaseMemObject(buffer1);
clReleaseMemObject(buffer2);
#else
TESTING_FREE_PIN( h_P );
TESTING_FREE_PIN( h_R );
#endif
for (j=0; j<tot_subs; j++) {
TESTING_FREE_DEV( d_lA[j] );
}
if (flag != 0)
break;
}
/* clean up */
for (i=0; i<num_gpus; i++) {
magma_queue_destroy( queues[2*i] );
magma_queue_destroy( queues[2*i+1] );
}
magma_finalize();
return 0;
}
int main( int argc, char** argv)
{
real_Double_t gpu_time, cpu_time;
double *h_A, *h_R, *VL, *VR, *h_work, *w1, *w2;
double *w1i, *w2i;
double c_neg_one = MAGMA_D_NEG_ONE;
double matnorm, tnrm, result[8];
/* Matrix size */
magma_int_t N=0, n2, lda, nb, lwork;
magma_int_t size[8] = {1024,2048,3072,4032,5184,6016,7040,8064};
magma_int_t i, j, info, checkres, once = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_vec_t jobl = MagmaVec;
magma_vec_t jobr = MagmaVec;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0) {
N = atoi(argv[++i]);
once = 1;
}
else if (strcmp("-LN", argv[i])==0)
jobl = MagmaNoVec;
else if (strcmp("-LV", argv[i])==0)
jobl = MagmaVec;
else if (strcmp("-RN", argv[i])==0)
jobr = MagmaNoVec;
else if (strcmp("-RV", argv[i])==0)
jobr = MagmaVec;
}
if ( N > 0 )
printf(" testing_dgeev -L[N|V] -R[N|V] -N %d\n\n", (int) N);
else
{
printf("\nUsage: \n");
printf(" testing_dgeev -L[N|V] -R[N|V] -N %d\n\n", 1024);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_dgeev -L[N|V] -R[N|V] -N %d\n\n", 1024);
N = size[7];
}
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
lda = N;
n2 = lda * N;
nb = magma_get_dgehrd_nb(N);
lwork = N*(2+nb);
// generous workspace - required by dget22
lwork = max(lwork, N * ( 5 + 2*N));
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( w1i, double, N );
TESTING_MALLOC_CPU( w2i, double, N );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_PIN( VL, double, n2 );
TESTING_MALLOC_PIN( VR, double, n2 );
TESTING_MALLOC_PIN( h_work, double, lwork );
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
printf(" N CPU Time(s) GPU Time(s) ||R||_F / ||A||_F\n");
printf("==========================================================\n");
for(i=0; i<8; i++){
if ( argc == 1 ){
N = size[i];
}
lda = N;
n2 = lda*N;
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// warm-up
magma_dgeev(jobl, jobr,
N, h_R, lda, w1, w1i,
VL, lda, VR, lda,
h_work, lwork, &info, queue);
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
gpu_time = magma_wtime();
magma_dgeev(jobl, jobr,
N, h_R, lda, w1, w1i,
VL, lda, VR, lda,
h_work, lwork, &info, queue);
gpu_time = magma_wtime() - gpu_time;
if (info < 0)
printf("Argument %d of magma_dgeev had an illegal value.\n", (int) -info);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dgeev(lapack_const(jobl), lapack_const(jobr),
&N, h_A, &lda, w2, w2i,
VL, &lda, VR, &lda,
h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info < 0)
printf("Argument %d of dgeev had an illegal value.\n", (int) -info);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
if ( checkres )
{
/* ===================================================================
Check the result following LAPACK's [zcds]drvev routine.
The following 7 tests are performed:
* (1) | A * VR - VR * W | / ( n |A| )
*
* Here VR is the matrix of unit right eigenvectors.
* W is a diagonal matrix with diagonal entries W(j).
*
* (2) | A**T * VL - VL * W**T | / ( n |A| )
*
* Here VL is the matrix of unit left eigenvectors, A**T is the
* ugate-transpose of A, and W is as above.
*
* (3) | |VR(i)| - 1 | and whether largest component real
*
* VR(i) denotes the i-th column of VR.
*
* (4) | |VL(i)| - 1 | and whether largest component real
*
* VL(i) denotes the i-th column of VL.
*
* (5) W(full) = W(partial)
*
* W(full) denotes the eigenvalues computed when both VR and VL
* are also computed, and W(partial) denotes the eigenvalues
* computed when only W, only W and VR, or only W and VL are
* computed.
*
* (6) VR(full) = VR(partial)
*
* VR(full) denotes the right eigenvectors computed when both VR
* and VL are computed, and VR(partial) denotes the result
* when only VR is computed.
*
* (7) VL(full) = VL(partial)
*
* VL(full) denotes the left eigenvectors computed when both VR
* and VL are also computed, and VL(partial) denotes the result
* when only VL is computed.
================================================================= */
int jj;
double ulp, ulpinv, vmx, vrmx, vtst, res[2];
double *LRE, DUM;
TESTING_MALLOC_PIN( LRE, double, n2 );
ulp = lapackf77_dlamch( "P" );
ulpinv = 1./ulp;
// Initialize RESULT
for (j = 0; j < 8; j++)
result[j] = -1.;
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dgeev(MagmaVec, MagmaVec,
N, h_R, lda, w1, w1i,
VL, lda, VR, lda,
h_work, lwork, &info, queue);
// Do test 1
lapackf77_dget22("N", "N", "N", &N, h_A, &lda, VR, &lda, w1, w1i,
h_work, res);
result[0] = res[0];
result[0] *= ulp;
// Do test 2
lapackf77_dget22("T", "N", "T", &N, h_A, &lda, VL, &lda, w1, w1i,
h_work, &result[1]);
result[1] *= ulp;
// Do test 3
result[2] = -1.;
for (j = 0; j < N; ++j) {
tnrm = 1.;
if (w1i[j] == 0.)
tnrm = cblas_dnrm2(N, &VR[j * lda], ione);
else if (w1i[j] > 0.)
tnrm = magma_dlapy2( cblas_dnrm2(N, &VR[j * lda], ione),
cblas_dnrm2(N, &VR[(j+1)* lda], ione) );
result[2] = fmax(result[2], fmin(ulpinv, magma_abs(tnrm-1.)/ulp));
if (w1i[j] > 0.)
{
vmx = vrmx = 0.;
for (jj = 0; jj <N; ++jj) {
vtst = magma_dlapy2( VR[jj+j*lda], VR[jj+(j+1)*lda]);
if (vtst > vmx)
vmx = vtst;
if ( (VR[jj + (j+1)*lda])==0. &&
magma_abs( VR[jj+j*lda] ) > vrmx)
vrmx = magma_abs( VR[jj+j*lda] );
}
if (vrmx / vmx < 1. - ulp * 2.)
result[2] = ulpinv;
}
}
result[2] *= ulp;
// Do test 4
result[3] = -1.;
for (j = 0; j < N; ++j) {
tnrm = 1.;
if (w1i[j] == 0.)
tnrm = cblas_dnrm2(N, &VL[j * lda], ione);
else if (w1i[j] > 0.)
tnrm = magma_dlapy2( cblas_dnrm2(N, &VL[j * lda], ione),
cblas_dnrm2(N, &VL[(j+1)* lda], ione) );
result[3] = fmax(result[3], fmin(ulpinv, magma_abs(tnrm-1.)/ulp));
if (w1i[j] > 0.)
{
vmx = vrmx = 0.;
for (jj = 0; jj <N; ++jj) {
vtst = magma_dlapy2( VL[jj+j*lda], VL[jj+(j+1)*lda]);
if (vtst > vmx)
vmx = vtst;
if ( (VL[jj + (j+1)*lda])==0. &&
magma_abs( VL[jj+j*lda]) > vrmx)
vrmx = magma_abs( VL[jj+j*lda] );
}
if (vrmx / vmx < 1. - ulp * 2.)
result[3] = ulpinv;
}
}
result[3] *= ulp;
// Compute eigenvalues only, and test them
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dgeev(MagmaNoVec, MagmaNoVec,
N, h_R, lda, w2, w2i,
&DUM, 1, &DUM, 1,
h_work, lwork, &info, queue);
if (info != 0) {
result[0] = ulpinv;
info = abs(info);
printf("Info = %d fo case N, N\n", (int) info);
}
// Do test 5
result[4] = 1;
for (j = 0; j < N; ++j)
if ( w1[j] != w2[j] || w1i[j] != w2i[j] )
result[4] = 0;
//if (result[4] == 0) printf("test 5 failed with N N\n");
// Compute eigenvalues and right eigenvectors, and test them
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dgeev(MagmaNoVec, MagmaVec,
N, h_R, lda, w2, w2i,
&DUM, 1, LRE, lda,
h_work, lwork, &info, queue);
if (info != 0) {
result[0] = ulpinv;
info = abs(info);
printf("Info = %d fo case N, V\n", (int) info);
}
// Do test 5 again
result[4] = 1;
for (j = 0; j < N; ++j)
if ( w1[j] != w2[j] || w1i[j] != w2i[j] )
result[4] = 0;
//if (result[4] == 0) printf("test 5 failed with N V\n");
// Do test 6
result[5] = 1;
for (j = 0; j < N; ++j)
for (jj = 0; jj < N; ++jj)
if ( VR[j+jj*lda] != LRE[j+jj*lda] )
result[5] = 0;
// Compute eigenvalues and left eigenvectors, and test them
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dgeev(MagmaVec, MagmaNoVec,
N, h_R, lda, w2, w2i,
LRE, lda, &DUM, 1,
h_work, lwork, &info, queue);
if (info != 0) {
result[0] = ulpinv;
info = abs(info);
printf("Info = %d fo case V, N\n", (int) info);
}
// Do test 5 again
result[4] = 1;
for (j = 0; j < N; ++j)
if ( w1[j] != w2[j] || w1i[j] != w2i[j] )
result[4] = 0;
//if (result[4] == 0) printf("test 5 failed with V N\n");
// Do test 7
result[6] = 1;
for (j = 0; j < N; ++j)
for (jj = 0; jj < N; ++jj)
if ( VL[j+jj*lda] != LRE[j+jj*lda] )
result[6] = 0;
printf("Test 1: | A * VR - VR * W | / ( n |A| ) = %e\n", result[0]);
printf("Test 2: | A'* VL - VL * W'| / ( n |A| ) = %e\n", result[1]);
printf("Test 3: | |VR(i)| - 1 | = %e\n", result[2]);
printf("Test 4: | |VL(i)| - 1 | = %e\n", result[3]);
printf("Test 5: W (full) == W (partial) = %f\n", result[4]);
printf("Test 6: VR (full) == VR (partial) = %f\n", result[5]);
printf("Test 7: VL (full) == VL (partial) = %f\n", result[6]);
//====================================================================
matnorm = lapackf77_dlange("f", &N, &ione, w1, &N, h_work);
blasf77_daxpy(&N, &c_neg_one, w1, &ione, w2, &ione);
result[7] = lapackf77_dlange("f", &N, &ione, w2, &N, h_work) / matnorm;
printf("%5d %6.2f %6.2f %e\n",
(int) N, cpu_time, gpu_time, result[7]);
TESTING_FREE_PIN( LRE );
}
else
{
printf("%5d %6.2f %6.2f\n",
(int) N, cpu_time, gpu_time);
}
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( w1i );
TESTING_FREE_CPU( w2i );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( VL );
TESTING_FREE_PIN( VR );
TESTING_FREE_PIN( h_work );
/* Shutdown */
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cswap, cswapblk, claswp, claswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaFloatComplex *h_A1, *h_A2;
magmaFloatComplex *h_R1, *h_R2;
magmaFloatComplex_ptr d_A1, d_A2;
// row-major and column-major performance
real_Double_t row_perf0 = MAGMA_D_NAN, col_perf0 = MAGMA_D_NAN;
real_Double_t row_perf1 = MAGMA_D_NAN, col_perf1 = MAGMA_D_NAN;
real_Double_t row_perf2 = MAGMA_D_NAN, col_perf2 = MAGMA_D_NAN;
real_Double_t row_perf4 = MAGMA_D_NAN;
real_Double_t row_perf5 = MAGMA_D_NAN, col_perf5 = MAGMA_D_NAN;
real_Double_t row_perf6 = MAGMA_D_NAN, col_perf6 = MAGMA_D_NAN;
real_Double_t row_perf7 = MAGMA_D_NAN;
real_Double_t cpu_perf = MAGMA_D_NAN;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magmaInt_ptr d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" %8s cswap cswap cswapblk claswp claswp2 claswpx ccopymatrix CPU (all in )\n", g_platform_str );
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj/col-maj row-blk/col-blk claswp (GByte/s)\n");
printf("=========================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_cgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaFloatComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
// getrf always makes ipiv[j] >= j+1, where ipiv is one based and j is zero based
// some implementations (e.g., MacOS dlaswp) assume this
for( j=0; j < nb; j++ ) {
ipiv[j] = (rand() % (N-j)) + j + 1;
assert( ipiv[j] >= j+1 );
assert( ipiv[j] <= N );
}
/* =====================================================================
* cublas / clBLAS / Xeon Phi cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasCswap( opts.handle, N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1 );
#else
magma_cswap( N, d_A1, ldda*j, 1, d_A2, ldda*(ipiv[j]-1), 1, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasCswap( opts.handle, N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
#else
magma_cswap( N, d_A1, j, ldda, d_A2, ipiv[j]-1, ldda, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswapblk, blocked version (2 matrices)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswp (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswp (1 matrix) - d_ipiv on GPU
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_claswp2( N, d_A1, ldda, 1, nb, d_ipiv, 1 );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswpx (extended for row- and col-major) (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
#endif
/* LAPACK swap on CPU for comparison */
time = magma_wtime();
lapackf77_claswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
#ifdef HAVE_CUBLAS
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_ccopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_ccopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf4, ((check & 0x040) != 0 ? '*' : ' '),
row_perf7, ((check & 0x080) != 0 ? '*' : ' '),
row_perf5, ((check & 0x100) != 0 ? '*' : ' '),
col_perf5, ((check & 0x200) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A, *h_R;
magma_int_t N, n2, lda, info;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu = %d, uplo = %s\n", (int) opts.ngpu, lapack_uplo_const(opts.uplo) );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
gflops = FLOPS_CPOTRF( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
magma_cmake_hpd( N, h_A, lda );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cpotrf( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_clange("f", &N, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double *h_A, *h_R;
double *d_A;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t N, n2, lda, ldda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" 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 ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_DPOTRI( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
magma_dmake_hpd( N, h_A, lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* factorize matrix */
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_dpotrf_gpu( opts.uplo, N, d_A, ldda, &info );
// check for exact singularity
//magma_dgetmatrix( N, N, d_A, ldda, h_R, lda );
//h_R[ 10 + 10*lda ] = MAGMA_D_MAKE( 0.0, 0.0 );
//magma_dsetmatrix( N, N, h_R, lda, d_A, ldda );
gpu_time = magma_wtime();
magma_dpotri_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dpotri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_dpotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_dpotri( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dpotri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_dgetmatrix( N, N, d_A, ldda, h_R, lda );
error = lapackf77_dlange("f", &N, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("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_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgetrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error;
double *h_A, *h_R;
magmaDouble_ptr d_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% M N CPU Gflop/s (ms) GPU Gflop/s (ms) Copy time (ms) ||PA-LU||/(||A||*N)\n");
printf("%%======================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = FLOPS_DGETRF( M, N ) / 1e9;
if ( N > 512 ) {
printf( "%5d %5d skipping because dgetf2 does not support N > 512\n", (int) M, (int) N );
continue;
}
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
real_Double_t set_time = magma_wtime();
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda, opts.queue );
set_time = magma_wtime() - set_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgetrf( &M, &N, h_A, &lda, ipiv, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_dgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( opts.queue );
magma_dgetf2_gpu( M, N, d_A, ldda, ipiv, opts.queue, &info );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_dgetf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
real_Double_t get_time = magma_wtime();
magma_dgetmatrix( M, N, d_A, ldda, h_A, lda, opts.queue );
get_time = magma_wtime() - get_time;
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f",
(int) M, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
set_time*1000.+get_time*1000.);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f",
(int) M, (int) N, gpu_perf, gpu_time*1000., set_time*1000.+get_time*1000. );
}
if ( opts.check ) {
magma_dgetmatrix( M, N, d_A, ldda, h_A, lda, opts.queue );
error = get_LU_error( M, N, h_R, lda, h_A, ipiv );
printf(" %8.2e %s\n", error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhesv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex *h_A, *h_B, *h_X, *work, temp;
real_Double_t gflops, gpu_perf, gpu_time = 0.0, cpu_perf=0, cpu_time=0;
double error, error_lapack = 0.0;
magma_int_t *ipiv;
magma_int_t N, n2, lda, ldb, sizeB, lwork, info;
magma_int_t status = 0, ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |Ax-b|/(N*|A|*|x|)\n");
printf("%%========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
ldb = N;
lda = N;
n2 = lda*N;
sizeB = ldb*opts.nrhs;
gflops = ( FLOPS_ZPOTRF( N ) + FLOPS_ZPOTRS( N, opts.nrhs ) ) / 1e9;
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_PIN( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_B, magmaDoubleComplex, sizeB );
TESTING_MALLOC_PIN( h_X, magmaDoubleComplex, sizeB );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lwork = -1;
lapackf77_zhesv(lapack_uplo_const(opts.uplo), &N, &opts.nrhs,
h_A, &lda, ipiv, h_X, &ldb, &temp, &lwork, &info);
lwork = (int)MAGMA_Z_REAL(temp);
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork );
init_matrix( N, N, h_A, lda );
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_zlacpy( MagmaFullStr, &N, &opts.nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_zhesv(lapack_uplo_const(opts.uplo), &N, &opts.nrhs,
h_A, &lda, ipiv, h_X, &ldb, work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_zhesv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
error_lapack = get_residual( opts.uplo, N, opts.nrhs, h_A, lda, ipiv, h_X, ldb, h_B, ldb );
TESTING_FREE_CPU( work );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( N, N, h_A, lda );
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_zlacpy( MagmaFullStr, &N, &opts.nrhs, h_B, &ldb, h_X, &ldb );
magma_setdevice(0);
gpu_time = magma_wtime();
magma_zhesv( opts.uplo, N, opts.nrhs, h_A, lda, ipiv, h_X, ldb, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_zhesv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) N, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) N, (int) N, gpu_perf, gpu_time );
}
if ( opts.check == 0 ) {
printf(" --- \n");
} else {
error = get_residual( opts.uplo, N, opts.nrhs, h_A, lda, ipiv, h_X, ldb, h_B, ldb );
printf(" %8.2e %s", error, (error < tol ? "ok" : "failed"));
if (opts.lapack)
printf(" (lapack rel.res. = %8.2e)", error_lapack);
printf("\n");
status += ! (error < tol);
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_PIN( h_X );
TESTING_FREE_PIN( h_B );
TESTING_FREE_PIN( h_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
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;
float tol;
opts.lapack |= (opts.version == 2 && opts.check == 2); // check (-c2) implies lapack (-l)
if ( opts.version != 2 && opts.check == 1 ) {
printf( "NOTE: version %d requires -c2 check due to the special structure of the\n"
"MAGMA cgeqrf results; using -c2.\n\n", (int) opts.version );
opts.check = 2;
}
printf( "version %d\n", (int) opts.version );
if ( opts.version == 2 ) {
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1*eps) ||I-Q'Q||_1 / (M*eps)\n");
printf("=========================================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
}
tol = 1.0;
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||Ax-b||_F/(N*||A||_F*||x||_F)\n");
printf("====================================================================================\n");
tol = opts.tolerance * lapackf77_slamch("E");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* Initialize the matrix */
for ( int j=0; j<4; j++ )
ISEED2[j] = ISEED[j]; // save seeds
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, 0, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
if ( opts.version == 2 ) {
magma_cgeqrf2_gpu( M, N, d_A, 0, ldda, tau, opts.queues2, &info );
}
else {
nb = magma_get_cgeqrf_nb( M );
size = (2*min(M, N) + (N+31)/32*32 )*nb;
TESTING_MALLOC_DEV( dT, magmaFloatComplex, size );
if ( opts.version == 1 ) {
magma_cgeqrf_gpu( M, N, d_A, 0, ldda, tau, dT, 0, opts.queue, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
magma_cgeqrf3_gpu( M, N, d_A, 0, ldda, tau, dT, opts.queue, &info );
}
#endif
else {
printf( "Unknown version %d\n", opts.version );
exit(1);
}
}
gpu_time = magma_wtime() - 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
=================================================================== */
magmaFloatComplex *tau2;
TESTING_MALLOC_CPU( tau2, magmaFloatComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau2, 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 ));
TESTING_FREE_CPU( tau2 );
}
if ( opts.check == 1 && M >= N ) {
/* =====================================================================
Check the result -- only version 1, cqrt02 requires M >= N
=================================================================== */
magma_int_t lwork = n2+N;
magmaFloatComplex *h_W1, *h_W2, *h_W3;
float *h_RW, results[2];
magma_cgetmatrix( M, N, d_A, 0, ldda, h_R, M, opts.queue );
TESTING_MALLOC_CPU( h_W1, magmaFloatComplex, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, magmaFloatComplex, n2 ); // R
TESTING_MALLOC_CPU( h_W3, magmaFloatComplex, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, float, M ); // RWORK
lapackf77_clarnv( &ione, ISEED2, &n2, h_A );
lapackf77_cqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
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] );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0], results[1] );
}
// todo also check results[1] < tol?
printf(" %s\n", (results[0] < tol ? "ok" : "failed"));
status += ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
}
else if ( opts.check == 2 && opts.version == 2 ) {
/* =====================================================================
Check the result compared to LAPACK -- only version 2
=================================================================== */
magma_cgetmatrix( M, N, d_A, 0, ldda, h_R, M, opts.queue );
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;
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 ? "ok" : "failed"));
status += ! (error < tol);
}
else if ( opts.check == 2 && M >= N ) {
/* =====================================================================
Check the result by solving linear system -- only versions 1 & 3, M >= N
=================================================================== */
magma_int_t lwork;
magmaFloatComplex *x, *b, *hwork;
magmaFloatComplex_ptr d_B;
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, magmaFloatComplex, N );
TESTING_MALLOC_CPU( b, magmaFloatComplex, M );
lapackf77_clarnv( &ione, ISEED, &N, x );
blasf77_cgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, M );
magma_csetvector( M, b, 1, d_B, 0, 1, opts.queue );
if ( opts.version == 1 ) {
// allocate hwork
magma_cgeqrs_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, tmp, -1, opts.queue, &info );
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, magmaFloatComplex, lwork );
// solve linear system
magma_cgeqrs_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, hwork, lwork, opts.queue, &info );
if (info != 0)
printf("magma_cgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_cgeqrs3_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, tmp, -1, opts.queue, &info );
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, magmaFloatComplex, lwork );
// solve linear system
magma_cgeqrs3_gpu( M, N, 1,
d_A, 0, ldda, tau, dT, 0,
d_B, 0, M, hwork, lwork, opts.queue, &info );
if (info != 0)
printf("magma_cgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", opts.version );
exit(1);
}
magma_cgetvector( N, d_B, 0, 1, x, 1, opts.queue );
// compute r = Ax - b, saved in b
lapackf77_clarnv( &ione, ISEED2, &n2, h_A );
blasf77_cgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
float norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_clange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_clange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_clange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (N * norm_A * norm_x);
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 ? "ok" : "failed"));
status += ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) ---",
(int) M, (int) N, gpu_perf, gpu_time);
}
printf("%s\n", (opts.check != 0 ? " (error check only for M >= N)" : ""));
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version != 2 )
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrs_gpu
*/
int main( int argc, char** argv)
{
//#if defined(PRECISION_s)
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double matnorm, work[1];
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_A2, *h_B, *h_X, *h_R, *tau, *hwork, tmp[1];
magmaDouble_ptr d_A, d_B;
/* Matrix size */
magma_int_t M = 0, N = 0, n2;
magma_int_t lda, ldb, ldda, lddb, lworkgpu, lhwork;
magma_int_t size[7] = {1024,2048,3072,4032,5184,6016,7000};
magma_int_t i, info, min_mn, nb, l1, l2;
magma_int_t ione = 1;
magma_int_t nrhs = 3;
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]);
else if (strcmp("-nrhs", argv[i])==0)
nrhs = atoi(argv[++i]);
}
if (N>0 && M>0 && M >= N)
printf(" testing_dgeqrs_gpu -nrhs %d -M %d -N %d\n\n", nrhs, M, N);
else
{
printf("\nUsage: \n");
printf(" testing_dgeqrs_gpu -nrhs %d -M %d -N %d\n\n", nrhs, M, N);
printf(" M has to be >= N, exit.\n");
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_dgeqrs_gpu -nrhs %d -M %d -N %d\n\n", nrhs, 1024, 1024);
M = N = size[6];
}
ldda = ((M+31)/32)*32;
lddb = ldda;
n2 = M * N;
min_mn = min(M, N);
nb = magma_get_dgeqrf_nb(M);
lda = ldb = M;
lworkgpu = (M-N + nb)*(nrhs+2*nb);
/* Allocate host memory for the matrix */
TESTING_MALLOC_PIN( tau, double, min_mn );
TESTING_MALLOC_PIN( h_A, double, lda*N );
TESTING_MALLOC_PIN( h_A2, double, lda*N );
TESTING_MALLOC_PIN( h_B, double, ldb*nrhs );
TESTING_MALLOC_PIN( h_X, double, ldb*nrhs );
TESTING_MALLOC_PIN( h_R, double, ldb*nrhs );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( d_B, double, lddb*nrhs );
/*
* Get size for host workspace
*/
lhwork = -1;
lapackf77_dgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
l1 = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lhwork = -1;
lapackf77_dormqr( MagmaLeftStr, MagmaTransStr,
&M, &nrhs, &min_mn, h_A, &lda, tau,
h_X, &ldb, tmp, &lhwork, &info);
l2 = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lhwork = max( max( l1, l2 ), lworkgpu );
TESTING_MALLOC_PIN( hwork, double, lhwork );
printf("\n");
printf(" ||b-Ax|| / (N||A||)\n");
printf(" M N CPU GFlop/s GPU GFlop/s CPU GPU \n");
printf("============================================================\n");
for(i=0; i<7; i++){
if (argc == 1){
M = N = size[i];
}
min_mn= min(M, N);
ldb = lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = (FLOPS_GEQRF( (double)M, (double)N )
+ FLOPS_GEQRS( (double)M, (double)N, (double)nrhs )) / 1e9;
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
n2 = M*nrhs;
lapackf77_dlarnv( &ione, ISEED, &n2, h_B );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* Warm up to measure the performance */
magma_dsetmatrix( M, N, h_A, 0, lda, d_A, 0, ldda, queue );
magma_dsetmatrix( M, nrhs, h_B, 0, ldb, d_B, 0, lddb, queue );
magma_dgels_gpu( MagmaNoTrans, M, N, nrhs, d_A, 0, ldda,
d_B, 0, lddb, hwork, lworkgpu, &info, queue);
magma_dsetmatrix( M, N, h_A, 0, lda, d_A, 0, ldda, queue );
magma_dsetmatrix( M, nrhs, h_B, 0, ldb, d_B, 0, lddb, queue );
gpu_time = magma_wtime();
magma_dgels_gpu( MagmaNoTrans, M, N, nrhs, d_A, 0, ldda,
d_B, 0, lddb, hwork, lworkgpu, &info, queue);
gpu_time = magma_wtime() - gpu_time;
if (info < 0)
printf("Argument %d of magma_dgels had an illegal value.\n", -info);
gpu_perf = gflops / gpu_time;
// Get the solution in h_X
magma_dgetmatrix( N, nrhs, d_B, 0, lddb, h_X, 0, ldb, queue );
// compute the residual
blasf77_dgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
matnorm = lapackf77_dlange("f", &M, &N, h_A, &lda, work);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_dgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, hwork, &lhwork, &info);
cpu_time = magma_wtime()-cpu_time;
cpu_perf = gflops / cpu_time;
if (info < 0)
printf("Argument %d of lapackf77_dgels had an illegal value.\n", -info);
blasf77_dgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb);
printf("%5d %5d %6.1f %6.1f %7.2e %7.2e\n",
M, N, cpu_perf, gpu_perf,
lapackf77_dlange("f", &M, &nrhs, h_B, &M, work)/(min_mn*matnorm),
lapackf77_dlange("f", &M, &nrhs, h_R, &M, work)/(min_mn*matnorm) );
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_B );
TESTING_FREE_PIN( h_X );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( hwork );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
/* Shutdown */
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error;
magmaFloatComplex *h_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |Ax-b|/(N*|A|*|x|)\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |PA-LU|/(N*|A|)\n");
}
printf("=========================================================================\n");
for( int 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_CGETRF( M, N ) / 1e9;
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_PIN( h_A, magmaFloatComplex, n2 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_cgetrf(&M, &N, h_A, &lda, ipiv, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( M, N, h_A, lda );
gpu_time = magma_wtime();
magma_cgetrf( M, N, h_A, lda, ipiv, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time );
}
if ( opts.check == 2 ) {
error = get_residual( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else if ( opts.check ) {
error = get_LU_error( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_PIN( h_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_zhemm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex calpha = MAGMA_Z_MAKE( 3.456, 5.678 );
magmaDoubleComplex cbeta = MAGMA_Z_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double error=0., errorbis=0., work[1];
magmaDoubleComplex *hA, *hX, *hB, *hR;
magmaDoubleComplex_ptr dA[MagmaMaxGPUs], dX[MagmaMaxGPUs], dB[MagmaMaxGPUs], dwork[MagmaMaxGPUs], hwork[MagmaMaxGPUs+1];
magmaDoubleComplex_ptr dA2;
magma_int_t M, N, size, lda, ldda, msize, nb, nstream;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
nstream = (opts.nstream > 0 ? opts.nstream : 2);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx = 0;
magma_buildconnection_mgpu(gnode, &nbcmplx, opts.ngpu);
printf("Initializing communication pattern... GPU-ncmplx %d\n\n", (int) nbcmplx);
for (int i=0; i < nbcmplx; ++i) {
int myngpu = gnode[i][MagmaMaxGPUs];
printf("cmplx %d has %d gpu ", i, myngpu);
for(int j=0; j < myngpu; ++j)
printf(" %d", (int) gnode[i][j]);
printf("\n");
}
magma_int_t nbevents = 2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t redevents[MagmaMaxGPUs][20];
magma_event_t redevents2[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
for( int d = 0; d < opts.ngpu; ++d ) {
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[d][i], cudaEventDisableTiming);
cudaEventCreateWithFlags(&redevents2[d][i], cudaEventDisableTiming);
}
}
printf( "nb %d, ngpu %d, nstream %d version %d\n", (int) nb, (int) opts.ngpu, (int) nstream, (int) opts.version );
printf(" M N nb offset CPU GFlop/s (sec) GPU GFlop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||X||\n");
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M;
ldda = ((M + 31)/32)*32;
size = lda*M;
gflops = FLOPS_ZHEMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = ldda*N*3;
magma_int_t hworksiz = lda*N;
TESTING_MALLOC_CPU( hA, magmaDoubleComplex, lda*M );
TESTING_MALLOC_CPU( hX, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( hB, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( hR, magmaDoubleComplex, lda*N );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( d );
TESTING_MALLOC_DEV( dA[d], magmaDoubleComplex, ldda*mlocal );
TESTING_MALLOC_DEV( dX[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dB[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork[d], magmaDoubleComplex, dworksiz );
TESTING_MALLOC_PIN( hwork[d], magmaDoubleComplex, hworksiz );
}
TESTING_MALLOC_PIN( hwork[opts.ngpu], magmaDoubleComplex, lda*N );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, magmaDoubleComplex, ldda*M );
}
lapackf77_zlarnv( &ione, iseed, &size, hA );
magma_zmake_hermitian( M, hA, lda );
size = lda*N;
lapackf77_zlarnv( &ione, iseed, &size, hX );
lapackf77_zlarnv( &ione, iseed, &size, hB );
lapackf77_zlacpy( "Full", &M, &N, hB, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
//magmablasSetKernelStream( streams[ d ][ 0 ] );
magma_zsetmatrix( M, N, hX, lda, dX[d], ldda );
//if (d == 0) magma_zsetmatrix( M, N, hB, lda, dB[d], ldda ); // this is wrong coz when offset != 0 the gpu who do the beta*C may be not 0 so this should be related to stdev(starting device who own i=0 first col)
magma_zsetmatrix( M, N, hB, lda, dB[d], ldda );
}
//memset(hR, 0, lda*N*sizeof(magmaDoubleComplex));
trace_init( 1, opts.ngpu, nstream, (magma_queue_t*) streams );
//magma_int_t offset = 0; //nb;
gpu_time = magma_sync_wtime(0);
magmablas_zhemm_mgpu_com(
MagmaLeft, MagmaLower, msize, N,
calpha, dA, ldda, offset,
dX, ldda,
cbeta, dB, ldda, dwork, dworksiz, hR, lda, hwork, hworksiz,
opts.ngpu, nb, streams, nstream, redevents2, nbevents, gnode, nbcmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "zhemm-m%d-n%d-nb%d-stream%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) nstream, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magmablasSetKernelStream( 0 );
magma_zsetmatrix( M, M, hA, lda, dA2, ldda );
magma_zsetmatrix( M, N, hX, lda, dX[0], ldda );
magma_zsetmatrix( M, N, hB, lda, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0);
magma_zhemm(
MagmaLeft, MagmaLower, msize, N,
calpha, dA2+offset*ldda+offset, ldda,
dX[0], ldda,
cbeta, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||X||
errorbis = lapackf77_zlange("fro", &msize, &msize, hA+offset*lda+offset, &lda, work );
errorbis *= lapackf77_zlange("fro", &msize, &N, hX, &lda, work );
//printf( "A =" ); magma_zprint( M, M, hA, lda );
//printf( "X =" ); magma_zprint( M, N, hX, lda );
//printf( "B =" ); magma_zprint( M, N, hB, lda );
cpu_time = magma_wtime();
blasf77_zhemm( "Left", "Lower", &msize, &N,
&calpha, hA+offset*lda+offset, &lda,
hX, &lda,
&cbeta, hB, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/*
trace_file = fopen("AJETE/C", "w");
for (int j = 0; j < N; j++)
for (int i = 0; i < siz; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, hB[j*lda+i]);
fclose(trace_file);
*/
magma_int_t firstprint=0;
for(magma_int_t dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_zgetmatrix( M, N, dB[dev], ldda, hR, lda );
// compute relative error ||R||/||A||*||X||, where R := B_magma - B_lapack = R - B
size = lda*N;
blasf77_zaxpy( &size, &c_neg_one, hB, &ione, hR, &ione );
error = lapackf77_zlange("fro", &msize, &N, hR, &lda, work) / errorbis;
//printf( "R =" ); magma_zprint( M, N, hR, lda );
if (firstprint == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (error < tol ? "ok" : "failed") );
}
else {
printf( "%89s %8.2e %s\n", " ",
error, (error < tol ? "ok" : "failed") );
}
status += ! (error < tol);
firstprint =1;
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hX );
TESTING_FREE_CPU( hB );
TESTING_FREE_PIN( hR );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
TESTING_FREE_DEV( dA[d] );
TESTING_FREE_DEV( dX[d] );
TESTING_FREE_DEV( dB[d] );
TESTING_FREE_DEV( dwork[d] );
TESTING_FREE_PIN( hwork[d] );
}
TESTING_FREE_PIN( hwork[opts.ngpu] );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( redevents[d][i] );
magma_event_destroy( redevents2[d][i] );
}
}
TESTING_FINALIZE();
return status;
}
// --------------------
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_time=0, cpu_perf=0, gpu_time, gpu_perf, mgpu_time, mgpu_perf, cuda_time, cuda_perf;
double Ynorm, error=0, error2=0, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t n_local[MagmaMaxGPUs];
magma_int_t N, Noffset, lda, ldda, blocks, lhwork, ldwork, matsize, vecsize;
magma_int_t incx = 1;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ylapack, *Ycublas, *Ymagma, *Ymagma1, *hwork;
magmaDoubleComplex_ptr dA, dX, dY;
magmaDoubleComplex_ptr d_lA[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magma_device_t dev;
magma_queue_t queues[MagmaMaxGPUs];
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code
double tol = opts.tolerance * lapackf77_dlamch("E");
magma_int_t nb = 64; // required by magmablas_zhemv_mgpu implementation
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_queue_create( dev, &queues[dev] );
}
// currently, tests all offsets in the offsets array;
// comment out loop below to test a specific offset.
magma_int_t offset = opts.offset;
magma_int_t offsets[] = { 0, 1, 31, 32, 33, 63, 64, 65, 100, 200 };
magma_int_t noffsets = sizeof(offsets) / sizeof(*offsets);
printf("%% uplo = %s, ngpu %d, block size = %d, offset %d\n",
lapack_uplo_const(opts.uplo), (int) opts.ngpu, (int) nb, (int) offset );
printf( "%% BLAS CUBLAS MAGMA 1 GPU MAGMA MGPU Error rel Error rel\n"
"%% N offset Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) to CUBLAS to LAPACK\n"
"%%==================================================================================================================\n" );
for( int itest = 0; itest < opts.ntest; ++itest ) {
// comment out these two lines & end of loop to test a specific offset
for( int ioffset=0; ioffset < noffsets; ioffset += 1 ) {
offset = offsets[ioffset];
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
Noffset = N + offset;
lda = Noffset;
ldda = magma_roundup( Noffset, opts.align ); // multiple of 32 by default
matsize = Noffset*ldda;
vecsize = (Noffset-1)*incx + 1;
gflops = FLOPS_ZHEMV( N ) / 1e9;
blocks = magma_ceildiv( N + (offset % nb), nb );
lhwork = N*opts.ngpu;
ldwork = ldda*(blocks + 1);
TESTING_MALLOC_CPU( A, magmaDoubleComplex, matsize );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ycublas, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma1, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ylapack, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( X, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( hwork, magmaDoubleComplex, lhwork );
magma_setdevice( opts.device );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, matsize );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, vecsize );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, vecsize );
// TODO make magma_zmalloc_bcyclic helper function?
for( dev=0; dev < opts.ngpu; dev++ ) {
n_local[dev] = ((Noffset/nb)/opts.ngpu)*nb;
if (dev < (Noffset/nb) % opts.ngpu)
n_local[dev] += nb;
else if (dev == (Noffset/nb) % opts.ngpu)
n_local[dev] += Noffset % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaDoubleComplex, ldda*n_local[dev] );
TESTING_MALLOC_DEV( dwork[dev], magmaDoubleComplex, ldwork );
}
//////////////////////////////////////////////////////////////////////////
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &matsize, A );
magma_zmake_hermitian( Noffset, A, lda );
lapackf77_zlarnv( &ione, ISEED, &vecsize, X );
lapackf77_zlarnv( &ione, ISEED, &vecsize, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_setdevice( opts.device );
magma_zsetmatrix( Noffset, Noffset, A, lda, dA, ldda, opts.queue );
magma_zsetvector( Noffset, X, incx, dX, incx, opts.queue );
magma_zsetvector( Noffset, Y, incx, dY, incx, opts.queue );
cuda_time = magma_sync_wtime(0);
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo), N,
&alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
&beta, dY + offset, incx );
cuda_time = magma_sync_wtime(0) - cuda_time;
cuda_perf = gflops / cuda_time;
magma_zgetvector( Noffset, dY, incx, Ycublas, incx, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (1 GPU)
=================================================================== */
magma_setdevice( opts.device );
magma_zsetvector( Noffset, Y, incx, dY, incx, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
magmablas_zhemv_work( opts.uplo, N,
alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
beta, dY + offset, incx, dwork[ opts.device ], ldwork,
opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_zgetvector( Noffset, dY, incx, Ymagma1, incx, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (multi-GPU)
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( Noffset, Noffset, A, lda, d_lA, ldda, opts.ngpu, nb, queues );
blasf77_zcopy( &Noffset, Y, &incx, Ymagma, &incx );
// workspaces do NOT need to be zero -- set to NAN to prove
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magmablas_zlaset( MagmaFull, ldwork, 1, MAGMA_Z_NAN, MAGMA_Z_NAN, dwork[dev], ldwork, opts.queue );
}
lapackf77_zlaset( "Full", &lhwork, &ione, &MAGMA_Z_NAN, &MAGMA_Z_NAN, hwork, &lhwork );
mgpu_time = magma_sync_wtime(0);
magma_int_t info;
info = magmablas_zhemv_mgpu(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if (info != 0) {
printf("magmablas_zhemv_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
info = magmablas_zhemv_mgpu_sync(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if (info != 0) {
printf("magmablas_zhemv_sync returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
mgpu_time = magma_sync_wtime(0) - mgpu_time;
mgpu_perf = gflops / mgpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
blasf77_zcopy( &Noffset, Y, &incx, Ylapack, &incx );
cpu_time = magma_wtime();
blasf77_zhemv( lapack_uplo_const(opts.uplo), &N,
&alpha, A + offset + offset*lda, &lda,
X + offset, &incx,
&beta, Ylapack + offset, &incx );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Compute the Difference LAPACK vs. Magma
=================================================================== */
Ynorm = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ylapack, &incx );
error2 = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work ) / Ynorm;
}
/* =====================================================================
Compute the Difference Cublas vs. Magma
=================================================================== */
Ynorm = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ycublas, &incx );
error = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work ) / Ynorm;
bool okay = (error < tol && error2 < tol);
status += ! okay;
if ( opts.lapack ) {
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) N, (int) offset,
cpu_perf, cpu_time*1000.,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, error2, (okay ? "ok" : "failed") );
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e --- %s\n",
(int) N, (int) offset,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, (okay ? "ok" : "failed") );
}
/* Free Memory */
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_CPU( Ymagma1 );
TESTING_FREE_CPU( Ylapack );
TESTING_FREE_PIN( X );
TESTING_FREE_PIN( hwork );
magma_setdevice( opts.device );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
for( dev=0; dev < opts.ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
// comment out these two lines line & top of loop test a specific offset
} // end for ioffset
printf( "\n" );
}
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_queue_destroy( queues[dev] );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time;
magmaDoubleComplex *h_A, *h_R, *h_B, *h_S, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
magma_int_t lrwork;
#endif
/* Matrix size */
double *w1, *w2, result[2]={0,0};
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_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
magma_range_t range = MagmaRangeAll;
if (opts.fraction != 1)
range = MagmaRangeI;
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: itype = %d, jobz = %s, range = %s, uplo = %s, opts.check = %d, fraction = %6.4f\n",
(int) opts.itype, lapack_vec_const(opts.jobz), lapack_range_const(range), lapack_uplo_const(opts.uplo),
(int) opts.check, opts.fraction);
printf(" N M GPU Time (sec)\n");
printf("============================\n");
magma_int_t threads = magma_get_parallel_numthreads();
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = magma_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;
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_S, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, double, lrwork);
#endif
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
magma_zmake_hpd( N, h_B, N );
magma_zmake_hermitian( N, h_A, N );
magma_int_t m1 = 0;
double vl = 0;
double vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (range == MagmaRangeI) {
il = 1;
iu = (int) (opts.fraction*N);
}
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if (opts.warmup) {
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_zhegvdx_2stage(opts.itype, opts.jobz, range, 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);
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_zhegvdx_2stage(opts.itype, opts.jobz, range, 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 ( opts.check && opts.jobz != MagmaNoVec ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | S(with V) - S(w/o V) | / | S |
=================================================================== */
#if defined(PRECISION_d) || defined(PRECISION_s)
double *rwork = h_work + N*N;
#endif
result[0] = 1.;
result[0] /= lapackf77_zlanhe("1", lapack_uplo_const(opts.uplo), &N, h_A, &N, rwork);
result[0] /= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_int_t m2 = m1;
lapackf77_zhegvd(&opts.itype, "N", lapack_uplo_const(opts.uplo), &N,
h_R, &N, h_S, &N, w2,
h_work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork,
&info);
double maxw=0, diff=0;
for(int j=0; j<m2; j++) {
maxw = max(maxw, fabs(w1[j]));
maxw = max(maxw, fabs(w2[j]));
diff = max(diff, fabs(w1[j] - w2[j]));
}
result[1] = diff / (m2*maxw);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %7.2f\n",
(int) N, (int) m1, gpu_time);
if ( opts.check && opts.jobz != MagmaNoVec ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype==1) {
printf(" | A Z - B Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
else if (opts.itype==2) {
printf(" | A B Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
else if (opts.itype==3) {
printf(" | B A Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
printf( " | D(w/ Z) - D(w/o Z) | / |D| = %8.2e %s\n\n", result[1], (result[1] < tolulp ? "ok" : "failed"));
status += ! (result[0] < tol && result[1] < tolulp);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_S );
TESTING_FREE_PIN( h_work );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_PIN( rwork );
#endif
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgelqf_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const float d_neg_one = MAGMA_D_NEG_ONE;
const float d_one = MAGMA_D_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float Anorm, error=0, error2=0;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex_ptr d_A;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |L - A*Q^H| |I - Q*Q^H|\n");
printf("%%==============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
n2 = lda*N;
nb = magma_get_cgeqrf_nb( M, N );
gflops = FLOPS_CGELQF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgelqf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max( lwork, M*nb );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_R, lda, d_A, ldda, opts.queue );
gpu_time = magma_wtime();
magma_cgelqf_gpu( M, N, d_A, ldda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cgelqf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result, following zlqt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
=================================================================== */
if ( opts.check ) {
magma_cgetmatrix( M, N, d_A, ldda, h_R, lda, opts.queue );
magma_int_t ldq = min_mn;
magma_int_t ldl = M;
magmaFloatComplex *Q, *L;
float *work;
TESTING_MALLOC_CPU( Q, magmaFloatComplex, ldq*N ); // K by N
TESTING_MALLOC_CPU( L, magmaFloatComplex, ldl*min_mn ); // M by K
TESTING_MALLOC_CPU( work, float, min_mn );
// generate K by N matrix Q, where K = min(M,N)
lapackf77_clacpy( "Upper", &min_mn, &N, h_R, &lda, Q, &ldq );
lapackf77_cunglq( &min_mn, &N, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy N by K matrix L
lapackf77_claset( "Upper", &M, &min_mn, &c_zero, &c_zero, L, &ldl );
lapackf77_clacpy( "Lower", &M, &min_mn, h_R, &lda, L, &ldl );
// error = || L - A*Q^H || / (N * ||A||)
blasf77_cgemm( "NoTrans", "Conj", &M, &min_mn, &N,
&c_neg_one, h_A, &lda, Q, &ldq, &c_one, L, &ldl );
Anorm = lapackf77_clange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_clange( "1", &M, &min_mn, L, &ldl, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set L = I (K by K), then L = I - Q*Q^H
// error = || I - Q*Q^H || / N
lapackf77_claset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, L, &ldl );
blasf77_cherk( "Upper", "NoTrans", &min_mn, &N, &d_neg_one, Q, &ldq, &d_one, L, &ldl );
error2 = safe_lapackf77_clanhe( "1", "Upper", &min_mn, L, &ldl, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( L ); L = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgelqf( &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_cgelqf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check ) {
bool okay = (error < tol && error2 < tol);
printf( "error %.4g, error2 %.4g, tol %.4g, okay %d\n", error, error2, tol, okay );
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetrf
*/
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time, error;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex_ptr d_A, dwork;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] = { 1024, 2048, 3072, 4032, 5184, 5600, 5600, 5600, 5600, 5600 };
magma_int_t ntest = 10;
magma_int_t i, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0, 0, 0, 1};
magmaDoubleComplex *work;
magmaDoubleComplex tmp;
double rwork[1];
magma_int_t *ipiv;
magma_int_t lwork, ldwork;
double A_norm, R_norm;
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[ntest-1] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_zgetri_gpu -N %d\n\n", 1024);
}
/* query for Lapack workspace size */
N = size[ntest-1];
lda = N;
work = &tmp;
lwork = -1;
lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
if (info != 0)
printf("lapackf77_zgetri returned error %d\n", (int) info);
lwork = int( MAGMA_Z_REAL( *work ));
/* query for Magma workspace size */
ldwork = N * magma_get_zgetri_nb( N );
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
/* Allocate memory */
n2 = N * N;
ldda = ((N+31)/32) * 32;
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork );
printf(" N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("========================================================\n");
for( i=0; i < ntest; i++ ){
N = size[i];
lda = N;
n2 = lda*N;
gflops = FLOPS_ZGETRI( (double)N ) / 1e9;
ldda = ((N+31)/32)*32;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
A_norm = lapackf77_zlange( "f", &N, &N, h_A, &lda, rwork );
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_zsetmatrix( N, N, h_A, 0, lda, d_A, 0, ldda, queue );
magma_zgetrf_gpu( N, N, d_A, 0, ldda, ipiv, &info, queue );
magma_zgetmatrix( N, N, d_A, 0, ldda, h_A, 0, lda, queue );
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
//warm-up
magma_zgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, &info, queue );
magma_zsetmatrix( N, N, h_A, 0, lda, d_A, 0, ldda, queue );
gpu_time = magma_wtime();
magma_zgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, &info, queue );
gpu_time = magma_wtime()-gpu_time;
if (info != 0)
printf("magma_zgetri_gpu returned error %d\n", (int) info);
gpu_perf = gflops / gpu_time;
magma_zgetmatrix( N, N, d_A, 0, ldda, h_R, 0, lda, queue );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_zgetri returned error %d\n", (int) info);
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_zaxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
R_norm = lapackf77_zlange( "f", &N, &N, h_R, &lda, rwork );
printf( "%5d %6.2f %6.2f %e\n",
(int) N, cpu_perf, gpu_perf, R_norm / A_norm );
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
/* Shutdown */
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
/* Constants */
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
/* Local variables */
real_Double_t gpu_time;
magmaFloatComplex *h_A, *h_R, *h_B, *h_S, *h_work;
#ifdef COMPLEX
float *rwork;
magma_int_t lrwork;
#endif
float *w1, *w2, result[2]={0,0};
magma_int_t *iwork;
magma_int_t N, n2, info, lda, lwork, liwork;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
magma_range_t range = MagmaRangeAll;
if (opts.fraction != 1)
range = MagmaRangeI;
// pass ngpu = -1 to test multi-GPU code using 1 gpu
magma_int_t abs_ngpu = abs( opts.ngpu );
printf("%% itype = %d, jobz = %s, range = %s, uplo = %s, fraction = %6.4f, ngpu = %d\n",
int(opts.itype), lapack_vec_const(opts.jobz), lapack_range_const(range), lapack_uplo_const(opts.uplo),
opts.fraction, int(abs_ngpu) );
if (opts.itype == 1) {
printf("%% N M GPU Time (sec) |AZ-BZD| |D - D_magma|\n");
}
else if (opts.itype == 2) {
printf("%% N M GPU Time (sec) |ABZ-ZD| |D - D_magma|\n");
}
else if (opts.itype == 3) {
printf("%% N M GPU Time (sec) |BAZ-ZD| |D - D_magma|\n");
}
printf("%%======================================================\n");
magma_int_t threads = magma_get_parallel_numthreads();
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
// TODO: test vl-vu range
magma_int_t m1 = 0;
float vl = 0;
float vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (opts.fraction == 0) {
il = max( 1, magma_int_t(0.1*N) );
iu = max( 1, magma_int_t(0.3*N) );
}
else {
il = 1;
iu = max( 1, magma_int_t(opts.fraction*N) );
}
magma_cheevdx_getworksize(N, threads, (opts.jobz == MagmaVec),
&lwork,
#ifdef COMPLEX
&lrwork,
#endif
&liwork);
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_S, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, max( lwork, N*N )); // check needs N*N
#ifdef COMPLEX
TESTING_MALLOC_PIN( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clarnv( &ione, ISEED, &n2, h_B );
magma_cmake_hpd( N, h_B, lda );
magma_cmake_hermitian( N, h_A, lda );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_B, &lda, h_S, &lda );
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
gpu_time = magma_wtime();
if (opts.ngpu == 1) {
magma_chegvdx_2stage( opts.itype, opts.jobz, range, opts.uplo,
N, h_R, lda, h_S, lda, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
else {
magma_chegvdx_2stage_m( abs_ngpu, opts.itype, opts.jobz, range, opts.uplo,
N, h_R, lda, h_S, lda, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork,
&info );
}
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_chegvdx_2stage returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A| |Z| N ) (itype = 1)
| A B Z - Z D | / ( |A| |Z| N ) (itype = 2)
| B A Z - Z D | / ( |A| |Z| N ) (itype = 3)
(2) | D(with V, magma) - D(w/o V, lapack) | / | D |
=================================================================== */
#ifdef REAL
float *rwork = h_work + N*N;
#endif
if ( opts.jobz != MagmaNoVec ) {
result[0] = 1.;
result[0] /= safe_lapackf77_clanhe("1", lapack_uplo_const(opts.uplo), &N, h_A, &lda, rwork);
result[0] /= lapackf77_clange("1", &N, &m1, h_R, &lda, rwork);
if (opts.itype == 1) {
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &lda, h_R, &lda, &c_zero, h_work, &N);
for (int i=0; i < m1; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_B, &lda, h_R, &lda, &c_one, h_work, &N);
result[0] *= lapackf77_clange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &lda, h_R, &lda, &c_zero, h_work, &N);
for (int i=0; i < m1; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_clange("1", &N, &m1, h_R, &lda, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &lda, h_R, &lda, &c_zero, h_work, &N);
for (int i=0; i < m1; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_clange("1", &N, &m1, h_R, &lda, rwork)/N;
}
}
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_B, &lda, h_S, &lda );
lapackf77_chegvd( &opts.itype, "N", lapack_uplo_const(opts.uplo), &N,
h_R, &lda, h_S, &lda, w2,
h_work, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork,
&info );
if (info != 0) {
printf("lapackf77_chegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
float maxw=0, diff=0;
for (int j=0; j < m1; j++) {
maxw = max(maxw, fabs(w1[j]));
maxw = max(maxw, fabs(w2[j]));
diff = max(diff, fabs(w1[j] - w2[j]));
}
result[1] = diff / (m1*maxw);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %9.4f ",
(int) N, (int) m1, gpu_time);
if ( opts.check ) {
bool okay = (result[1] < tolulp);
if ( opts.jobz != MagmaNoVec ) {
okay = okay && (result[0] < tol);
printf(" %8.2e", result[0] );
}
else {
printf(" --- ");
}
printf(" %8.2e %s\n", result[1], (okay ? "ok" : "failed"));
status += ! okay;
}
else {
printf(" ---\n");
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_S );
TESTING_FREE_PIN( h_work );
#ifdef COMPLEX
TESTING_FREE_PIN( rwork );
#endif
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_T, *h_R, *tau, *h_work, tmp[1];
magmaDouble_ptr d_A, d_T, ddA, dtau;
magmaDouble_ptr dwork;
/* Matrix size */
magma_int_t M = 0, N = 0, n2, lda, ldda, lwork;
const int MAXTESTS = 10;
magma_int_t msize[MAXTESTS] = { 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 8100, 8192 };
magma_int_t nsize[MAXTESTS] = { 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 8100, 8192 };
magma_int_t i, info, min_mn;
magma_int_t ione = 1;
//magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t checkres;
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
// process command line arguments
printf( "\nUsage: %s -N <m,n> -c\n", argv[0] );
printf( " -N can be repeated up to %d times. If only m is given, then m=n.\n", MAXTESTS );
printf( " -c or setting $MAGMA_TESTINGS_CHECK runs LAPACK and checks result.\n\n" );
int ntest = 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 );
int m, n;
info = sscanf( argv[++i], "%d,%d", &m, &n );
if ( info == 2 && m > 0 && n > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = n;
}
else if ( info == 1 && m > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = m; // implicitly
}
else {
printf( "error: -N %s is invalid; ensure m > 0, n > 0.\n", argv[i] );
exit(1);
}
M = max( M, msize[ ntest ] );
N = max( N, nsize[ ntest ] );
ntest++;
}
else if ( strcmp("-M", argv[i]) == 0 ) {
printf( "-M has been replaced in favor of -N m,n to allow -N to be repeated.\n\n" );
exit(1);
}
else if ( strcmp("-c", argv[i]) == 0 ) {
checkres = true;
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( ntest == 0 ) {
ntest = MAXTESTS;
M = msize[ntest-1];
N = nsize[ntest-1];
}
ldda = ((M+31)/32)*32;
n2 = M * N;
min_mn = min(M, N);
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
/* Allocate memory for the matrix */
TESTING_MALLOC_PIN( tau, double, min_mn );
TESTING_MALLOC_PIN( h_A, double, n2 );
TESTING_MALLOC_PIN( h_T, double, N*N );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( d_T, double, N*N );
TESTING_MALLOC_DEV( ddA, double, N*N );
TESTING_MALLOC_DEV( dtau, double, min_mn );
TESTING_MALLOC_DEV( dwork, double, max(5*min_mn, (32*2+2)*min_mn) );
double *h1 = (double*)malloc(sizeof(double)*N*N);
memset(h1, 0, N*N*sizeof(double));
clEnqueueWriteBuffer(queue, ddA, CL_TRUE, 0, sizeof(double)*N*N, h1, 0, NULL, NULL);
clEnqueueWriteBuffer(queue, d_T, CL_TRUE, 0, sizeof(double)*N*N, h1, 0, NULL, NULL);
lwork = -1;
lapackf77_dgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_PIN( h_work, double, lwork );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F/||A||_F ||R_T||\n");
printf("=============================================================================\n");
for( i = 0; i < ntest; ++i ) {
M = msize[i];
N = nsize[i];
min_mn= min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = (FLOPS_DGEQRF( M, N ) + FLOPS_DGEQRT( M, N)) / 1e9;
/* Initialize the matrix */
magma_int_t ISEED[4] = {0,0,0,1};
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( M, N, h_R, 0, lda, d_A, 0, ldda, queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// warm-up
// magma_dgeqr2x3_gpu(&M, &N, d_A, 0, &ldda, dtau, 0, d_T, 0, ddA, 0, dwork, 0, &info, queue);
/*
magma_dsetmatrix( M, N, h_R, 0, lda, d_A, 0, ldda, queue );
clEnqueueWriteBuffer(queue, ddA, CL_TRUE, 0, sizeof(double)*N*N, h1, 0, NULL, NULL);
clEnqueueWriteBuffer(queue, d_T, CL_TRUE, 0, sizeof(double)*N*N, h1, 0, NULL, NULL);
*/
gpu_time = magma_wtime();
magma_dgeqr2x3_gpu(&M, &N, d_A, 0, &ldda, dtau, 0, d_T, 0, ddA, 0, dwork, 0, &info, queue);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgeqrf returned error %d.\n", (int) info);
if ( checkres ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&M, &N, h_A, &lda, tau, h_work, &N);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgeqrf returned error %d.\n", (int) info);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_dgetmatrix( M, N, d_A, 0, ldda, h_R, 0, M, queue );
magma_dgetmatrix( N, N, ddA, 0, N, h_T, 0, N, queue );
// Restore the upper triangular part of A before the check
for(int col=0; col<N; col++){
for(int row=0; row<=col; row++)
h_R[row + col*M] = h_T[row + col*N];
}
error = lapackf77_dlange("M", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("M", &M, &N, h_R, &lda, work) / error;
// Check if T is the same
double terr = 0.;
magma_dgetmatrix( N, N, d_T, 0, N, h_T, 0, N, queue );
for(int col=0; col<N; col++)
for(int row=0; row<=col; row++)
terr += ( MAGMA_D_ABS(h_work[row + col*N] - h_T[row + col*N])*
MAGMA_D_ABS(h_work[row + col*N] - h_T[row + col*N]) );
terr = magma_dsqrt(terr);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, terr);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time);
}
}
/* Memory clean up */
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_T );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_T );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( dtau );
free(h1);
magma_queue_destroy( queue );
magma_finalize();
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
float matnorm, work[1];
float mzone = MAGMA_S_NEG_ONE;
float *h_A, *h_R, *tau, *hwork, tmp[1];
magmaFloat_ptr d_A;
/* Matrix size */
magma_int_t M = 0, N = 0, n2, lda, ldda, lhwork;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10176};
magma_int_t i, info, min_mn;
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 (M>0 && N>0)
printf(" testing_sgeqrf_gpu -M %d -N %d\n\n", M, N);
else
{
printf("\nUsage: \n");
printf(" testing_sgeqrf_gpu -M %d -N %d\n\n", 1024, 1024);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_sgeqrf_gpu -M %d -N %d\n\n", 1024, 1024);
M = N = size[7];
}
/* Initialize */
magma_queue_t queue1, queue2;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue1 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue2 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
magma_queue_t queues[2] = {queue1, queue2};
ldda = ((M+31)/32)*32;
n2 = M * N;
min_mn = min(M, N);
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
lhwork = -1;
lapackf77_sgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
lhwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lhwork );
printf("\n\n");
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("======================================================================\n");
for(i=0; i<8; i++){
if (argc == 1){
M = N = size[i];
}
min_mn= min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS( (float)M, (float)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sgeqrf(&M, &N, h_A, &M, tau, hwork, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info < 0)
printf("Argument %d of lapack_sgeqrf had an illegal value.\n", -info);
cpu_perf = gflops / cpu_time;
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix( M, N, h_R, 0, lda, d_A, 0, ldda, queue1 );
magma_sgeqrf2_gpu( M, N, d_A, 0, ldda, tau, &info, queues);
magma_ssetmatrix( M, N, h_R, 0, lda, d_A, 0, ldda, queue1 );
clFinish(queue1);
clFinish(queue2);
gpu_time = magma_wtime();
magma_sgeqrf2_gpu( M, N, d_A, 0, ldda, tau, &info, queues);
gpu_time = magma_wtime() - gpu_time;
if (info < 0)
printf("Argument %d of magma_sgeqrf2 had an illegal value.\n", -info);
gpu_perf = gflops / gpu_time;
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_sgetmatrix( M, N, d_A, 0, ldda, h_R, 0, M, queue1 );
matnorm = lapackf77_slange("f", &M, &N, h_A, &M, work);
blasf77_saxpy(&n2, &mzone, h_A, &ione, h_R, &ione);
printf("%5d %5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
M, N, cpu_perf, cpu_time, gpu_perf, gpu_time,
lapackf77_slange("f", &M, &N, h_R, &M, work) / matnorm);
if (argc != 1)
break;
}
/* clean up */
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( hwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
magma_queue_destroy( queue1 );
magma_queue_destroy( queue2 );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetrf
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R, *work;
magmaDoubleComplex_ptr d_A, dwork;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, ldda, info, lwork, ldwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex tmp;
double error, rwork[1];
magma_int_t *ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// need looser bound (3000*eps instead of 30*eps) for tests
// TODO: should compute ||I - A*A^{-1}|| / (n*||A||*||A^{-1}||)
opts.tolerance = max( 3000., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / (N*||A||_F)\n");
printf("=================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
ldwork = N * magma_get_zgetri_nb( N );
gflops = FLOPS_ZGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0)
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = int( MAGMA_Z_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
error = lapackf77_zlange( "f", &N, &N, h_A, &lda, rwork ); // norm(A)
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_zsetmatrix( N, N, h_A, lda, d_A, 0, ldda, opts.queue );
magma_zgetrf_gpu( N, N, d_A, 0, ldda, ipiv, opts.queue, &info );
magma_zgetmatrix( N, N, d_A, 0, ldda, h_A, lda, opts.queue );
if ( info != 0 )
printf("magma_zgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_A, lda, d_A, 0, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, opts.queues2, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_zgetmatrix( N, N, d_A, 0, ldda, h_R, lda, opts.queue );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_zaxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_zlange( "f", &N, &N, h_R, &lda, rwork ) / (N*error);
printf( "%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf( "%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
real_Double_t mgpu_time;
float *h_A, *h_Ainit, *h_B, *h_Binit, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork;
#endif
float *w1, result=0;
magma_int_t *iwork;
magma_int_t N, n2, info, lwork, liwork;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
magma_range_t range = MagmaRangeAll;
if (opts.fraction != 1)
range = MagmaRangeI;
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: ngpu = %d, itype = %d, jobz = %s, range = %s, uplo = %s, opts.check = %d, fraction = %6.4f\n",
(int) opts.ngpu, (int) opts.itype,
lapack_vec_const(opts.jobz), lapack_range_const(range), lapack_uplo_const(opts.uplo),
(int) opts.check, opts.fraction);
printf(" N M ngpu MGPU Time (sec)\n");
printf("====================================\n");
magma_int_t threads = magma_get_parallel_numthreads();
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = magma_sbulge_get_lq2(N, threads) + 2*N + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = magma_sbulge_get_lq2(N, threads) + 1 + 6*N + 2*N*N;
#endif
liwork = 3 + 5*N;
//magma_int_t NB = 96;//magma_bulge_get_nb(N);
//magma_int_t sizvblg = magma_sbulge_get_lq2(N, threads);
//magma_int_t siz = max(sizvblg,n2)+2*(N*NB+N)+24*N;
/* Allocate host memory for the matrix */
TESTING_MALLOC_PIN( h_A, float, n2 );
TESTING_MALLOC_PIN( h_B, float, n2 );
TESTING_MALLOC_PIN( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, float, lrwork);
#endif
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork);
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
magma_smake_hpd( N, h_B, N );
magma_smake_symmetric( N, h_A, N );
if ( opts.warmup || opts.check ) {
TESTING_MALLOC_CPU( h_Ainit, float, n2 );
TESTING_MALLOC_CPU( h_Binit, float, n2 );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_Binit, &N );
}
magma_int_t m1 = 0;
float vl = 0;
float vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (range == MagmaRangeI) {
il = 1;
iu = (int) (opts.fraction*N);
}
if ( opts.warmup ) {
// ==================================================================
// Warmup using MAGMA. I prefer to use smalltest to warmup A-
// ==================================================================
magma_ssygvdx_2stage_m(opts.ngpu, opts.itype, opts.jobz, range, opts.uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
mgpu_time = magma_wtime();
magma_ssygvdx_2stage_m(opts.ngpu, opts.itype, opts.jobz, range, opts.uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
mgpu_time = magma_wtime() - mgpu_time;
if ( opts.check ) {
// ===================================================================
// Check the results following the LAPACK's [zc]hegvdx routine.
// A x = lambda B x is solved
// and the following 3 tests computed:
// (1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
// | A B Z - Z D | / ( |A||Z| N ) (itype = 2)
// | B A Z - Z D | / ( |A||Z| N ) (itype = 3)
// ===================================================================
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = h_work + N*N;
#endif
result = 1.;
result /= lapackf77_slansy("1", lapack_uplo_const(opts.uplo), &N, h_Ainit, &N, rwork);
result /= lapackf77_slange("1", &N , &m1, h_A, &N, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result *= lapackf77_slange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Binit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_slange("1", &N, &m1, h_A, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_sscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_ssymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_slange("1", &N, &m1, h_A, &N, rwork)/N;
}
}
// ===================================================================
// Print execution time
// ===================================================================
printf("%5d %5d %4d %7.2f\n",
(int) N, (int) m1, (int) opts.ngpu, mgpu_time);
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype==1) {
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e %s\n", result, (result < tol ? "ok" : "failed") );
}
else if (opts.itype==2) {
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result, (result < tol ? "ok" : "failed") );
}
else if (opts.itype==3) {
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result, (result < tol ? "ok" : "failed") );
}
printf("\n");
status += ! (result < tol);
}
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_B );
TESTING_FREE_PIN( h_work );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_PIN( rwork );
#endif
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( iwork );
if ( opts.warmup || opts.check ) {
TESTING_FREE_CPU( h_Ainit );
TESTING_FREE_CPU( h_Binit );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE_MGPU();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_T, *h_R, *tau, *h_work, tmp[1];
double *d_A, *d_T, *ddA, *dtau;
double *d_A2, *d_T2, *ddA2, *dtau2;
double *dwork, *dwork2;
magma_int_t M, N, lda, ldda, lwork, n2, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
#define BLOCK_SIZE 64
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = 10. * opts.tolerance * lapackf77_dlamch("E");
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
printf("version %d\n", (int) opts.version );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F/||A||_F ||R_T||\n");
printf("=============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if (N > 128) {
printf("%5d %5d skipping because dgeqr2x requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because dgeqr2x requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = (FLOPS_DGEQRF( M, N ) + FLOPS_DGEQRT( M, N )) / 1e9;
/* Allocate memory for the matrix */
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_T, double, N*N );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( d_T, double, N*N );
TESTING_MALLOC_DEV( ddA, double, N*N );
TESTING_MALLOC_DEV( dtau, double, min_mn );
TESTING_MALLOC_DEV( d_A2, double, ldda*N );
TESTING_MALLOC_DEV( d_T2, double, N*N );
TESTING_MALLOC_DEV( ddA2, double, N*N );
TESTING_MALLOC_DEV( dtau2, double, min_mn );
TESTING_MALLOC_DEV( dwork, double, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
TESTING_MALLOC_DEV( dwork2, double, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
// todo replace with magma_dlaset
cudaMemset(ddA, 0, N*N*sizeof(double));
cudaMemset(d_T, 0, N*N*sizeof(double));
cudaMemset(ddA2, 0, N*N*sizeof(double));
cudaMemset(d_T2, 0, N*N*sizeof(double));
lwork = -1;
lapackf77_dgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_CPU( h_work, double, lwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda );
magma_dsetmatrix( M, N, h_R, lda, d_A2, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime(0);
if (opts.version == 1)
magma_dgeqr2x_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 2)
magma_dgeqr2x2_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 3)
magma_dgeqr2x3_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else {
printf( "call magma_dgeqr2x4_gpu\n" );
/*
Going through NULL stream is faster
Going through any stream is slower
Doing two streams in parallel is slower than doing them sequentially
Queuing happens on the NULL stream - user defined buffers are smaller?
*/
magma_dgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, NULL);
//magma_dgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, stream[1]);
//magma_dgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, stream[0]);
//magma_dgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, NULL);
//gflops *= 2;
}
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_dgeqr2x_gpu version %d returned error %d: %s.\n",
(int) opts.version, (int) info, magma_strerror( info ));
}
else {
if ( opts.check ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&M, &N, h_A, &lda, tau, h_work, &N);
//magma_dgeqr2(&M, &N, h_A, &lda, tau, h_work, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_dgetmatrix( M, N, d_A, ldda, h_R, M );
magma_dgetmatrix( N, N, ddA, N, h_T, N );
// Restore the upper triangular part of A before the check
for(int col=0; col < N; col++){
for(int row=0; row <= col; row++)
h_R[row + col*M] = h_T[row + col*N];
}
error = lapackf77_dlange("M", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("M", &M, &N, h_R, &lda, work) / (N * error);
// Check if T is the same
magma_dgetmatrix( N, N, d_T, N, h_T, N );
double terr = 0.;
for(int col=0; col < N; col++)
for(int row=0; row <= col; row++)
terr += ( MAGMA_D_ABS(h_work[row + col*N] - h_T[row + col*N])*
MAGMA_D_ABS(h_work[row + col*N] - h_T[row + col*N]) );
terr = magma_dsqrt(terr);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, terr, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time);
}
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_T );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_T );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( d_A2 );
TESTING_FREE_DEV( d_T2 );
TESTING_FREE_DEV( ddA2 );
TESTING_FREE_DEV( dtau2 );
TESTING_FREE_DEV( dwork2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zungqr_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double Anorm, error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *hA, *hR, *tau, *h_work;
magmaDoubleComplex_ptr dA, dT;
magma_int_t m, n, k;
magma_int_t n2, lda, ldda, lwork, min_mn, nb, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" m n k CPU GFlop/s (sec) GPU GFlop/s (sec) ||R|| / ||A||\n");
printf("=========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
if ( m < n || n < k ) {
printf( "%5d %5d %5d skipping because m < n or n < k\n", (int) m, (int) n, (int) k );
continue;
}
lda = m;
ldda = ((m + 31)/32)*32;
n2 = lda*n;
min_mn = min(m, n);
nb = magma_get_zgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_ZUNGQR( m, n, k ) / 1e9;
TESTING_MALLOC_PIN( hA, magmaDoubleComplex, lda*n );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( hR, magmaDoubleComplex, lda*n );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*n );
TESTING_MALLOC_DEV( dT, magmaDoubleComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_zlarnv( &ione, ISEED, &n2, hA );
lapackf77_zlacpy( MagmaFullStr, &m, &n, hA, &lda, hR, &lda );
Anorm = lapackf77_zlange("f", &m, &n, hA, &lda, work );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// first, get QR factors in both hA and dA
// okay that magma_zgeqrf_gpu has special structure for R; R isn't used here.
magma_zsetmatrix( m, n, hA, lda, dA, ldda );
magma_zgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_zgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_zgetmatrix( m, n, dA, ldda, hA, lda );
gpu_time = magma_wtime();
magma_zungqr_gpu( m, n, k, dA, ldda, tau, dT, nb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get dA back to the CPU to compare with the CPU result.
magma_zgetmatrix( m, n, dA, ldda, hR, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zungqr( &m, &n, &k, hA, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_zaxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_zlange("f", &m, &n, hR, &lda, work) / Anorm;
bool okay = (error < tol);
status += ! okay;
printf("%5d %5d %5d %7.1f (%7.2f) %7.1f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (okay ? "ok" : "failed"));
}
else {
printf("%5d %5d %5d --- ( --- ) %7.1f (%7.2f) --- \n",
(int) m, (int) n, (int) k,
gpu_perf, gpu_time );
}
TESTING_FREE_PIN( hA );
TESTING_FREE_PIN( h_work );
TESTING_FREE_CPU( hR );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}