/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgemm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
double magma_error, cublas_error, Cnorm, work[1];
magma_int_t M, N, K;
magma_int_t Am, An, Bm, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
double *h_A, *h_B, *h_C, *h_Cmagma, *h_Ccublas;
double *d_A, *d_B, *d_C;
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
double beta = MAGMA_D_MAKE( -0.48, 0.38 );
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n");
printf("transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf(" M N K MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("=========================================================================================================\n");
for( int 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];
gflops = FLOPS_DGEMM( M, N, K ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if ( opts.transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
ldc = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*An;
sizeB = ldb*Bn;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, double, lda*An );
TESTING_MALLOC_CPU( h_B, double, ldb*Bn );
TESTING_MALLOC_CPU( h_C, double, ldc*N );
TESTING_MALLOC_CPU( h_Cmagma, double, ldc*N );
TESTING_MALLOC_CPU( h_Ccublas, double, ldc*N );
TESTING_MALLOC_DEV( d_A, double, ldda*An );
TESTING_MALLOC_DEV( d_B, double, lddb*Bn );
TESTING_MALLOC_DEV( d_C, double, lddc*N );
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_dlarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_dlarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_dsetmatrix( Am, An, h_A, lda, d_A, ldda );
magma_dsetmatrix( Bm, Bn, h_B, ldb, d_B, lddb );
magma_dsetmatrix( M, N, h_C, ldc, d_C, lddc );
magma_time = magma_sync_wtime( NULL );
magmablas_dgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, ldda,
d_B, lddb,
beta, d_C, lddc );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_dgetmatrix( M, N, d_C, lddc, h_Cmagma, ldc );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_dsetmatrix( M, N, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasDgemm( handle, cublas_trans_const(opts.transA), cublas_trans_const(opts.transB), M, N, K,
&alpha, d_A, ldda,
d_B, lddb,
&beta, d_C, lddc );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_dgetmatrix( M, N, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_dgemm( lapack_trans_const(opts.transA), lapack_trans_const(opts.transB), &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_dlange( "M", &M, &N, h_C, &ldc, work );
blasf77_daxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cmagma, &ione );
magma_error = lapackf77_dlange( "M", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
blasf77_daxpy( &sizeC, &c_neg_one, h_C, &ione, h_Ccublas, &ione );
cublas_error = lapackf77_dlange( "M", &M, &N, h_Ccublas, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
else {
// compute relative error for magma, relative to cublas
Cnorm = lapackf77_dlange( "M", &M, &N, h_Ccublas, &ldc, work );
blasf77_daxpy( &sizeC, &c_neg_one, h_Ccublas, &ione, h_Cmagma, &ione );
magma_error = lapackf77_dlange( "M", &M, &N, h_Cmagma, &ldc, work ); // / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e --- %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error,
(magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Cmagma );
TESTING_FREE_CPU( h_Ccublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dlat2s and slat2d
*/
int main( int argc, char** argv )
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define SA(i_,j_) (SA + (i_) + (j_)*lda)
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
float serror, swork[1];
double c_neg_one = MAGMA_D_NEG_ONE;
float s_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t n, lda, ldda, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
float *SA, *SR;
double *A, *R;
float *dSA;
double *dA;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper };
printf("func uplo N CPU GB/s (ms) GPU GB/s (ms) ||R||_F\n");
printf("=====================================================================\n");
for( int iuplo = 0; iuplo < 2; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
n = opts.nsize[itest];
lda = n;
ldda = ((n+31)/32)*32;
// 0.5*(n+1)*n double-real loads and 0.5*(n+1)*n single-real stores (and vice-versa for slat2d)
gbytes = (real_Double_t) 0.5*(n+1)*n * (sizeof(double) + sizeof(float)) / 1e9;
size = ldda*n; // ldda >= lda
TESTING_MALLOC_CPU( SA, float, size );
TESTING_MALLOC_CPU( A, double, size );
TESTING_MALLOC_CPU( SR, float, size );
TESTING_MALLOC_CPU( R, double, size );
TESTING_MALLOC_DEV( dSA, float, size );
TESTING_MALLOC_DEV( dA, double, size );
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_slarnv( &ione, ISEED, &size, SA );
magma_dsetmatrix( n, n, A, lda, dA, ldda );
magma_ssetmatrix( n, n, SA, lda, dSA, ldda );
/* =====================================================================
Performs operation using LAPACK dlat2s
=================================================================== */
info = 0;
cpu_time = magma_wtime();
lapackf77_dlat2s( lapack_uplo_const(uplo[iuplo]), &n, A, &lda, SA, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if (info != 0)
printf("lapackf77_dlat2s returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA dlat2s
=================================================================== */
gpu_time = magma_sync_wtime(0);
magmablas_dlat2s( uplo[iuplo], n, dA, ldda, dSA, ldda, &info );
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gbytes / gpu_time;
if (info != 0)
printf("magmablas_dlat2s returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_sgetmatrix( n, n, dSA, ldda, SR, lda );
if ( opts.verbose ) {
printf( "A= " ); magma_dprint( n, n, A, lda );
printf( "SA= " ); magma_sprint( n, n, SA, lda );
printf( "dA= " ); magma_dprint_gpu( n, n, dA, ldda );
printf( "dSA=" ); magma_sprint_gpu( n, n, dSA, ldda );
}
/* =====================================================================
compute error |SA_magma - SA_lapack|
should be zero if both are IEEE compliant
=================================================================== */
blasf77_saxpy( &size, &s_neg_one, SA, &ione, SR, &ione );
serror = lapackf77_slange( "Fro", &n, &n, SR, &lda, swork );
printf( "dlat2s %5s %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
lapack_uplo_const(uplo[iuplo]), (int) n,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
serror, (serror == 0 ? "ok" : "failed") );
status += ! (serror == 0);
/* =====================================================================
Reset matrices
=================================================================== */
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_slarnv( &ione, ISEED, &size, SA );
magma_dsetmatrix( n, n, A, lda, dA, ldda );
magma_ssetmatrix( n, n, SA, lda, dSA, ldda );
/* =====================================================================
Performs operation using LAPACK slat2d
LAPACK doesn't implement slat2d; use our own simple implementation.
=================================================================== */
cpu_time = magma_wtime();
if ( uplo[iuplo] == MagmaLower ) {
for( int j=0; j < n; ++j ) {
for( int i=j; i < n; ++i ) {
*A(i,j) = MAGMA_D_MAKE( real(*SA(i,j)), imag(*SA(i,j)) );
}
}
}
else { // upper
for( int j=0; j < n; ++j ) {
for( int i=0; i <= j; ++i ) {
*A(i,j) = MAGMA_D_MAKE( real(*SA(i,j)), imag(*SA(i,j)) );
}
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if (info != 0)
printf("lapackf77_slat2d returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA slat2d
=================================================================== */
magma_ssetmatrix( n, n, SA, lda, dSA, ldda );
gpu_time = magma_sync_wtime(0);
magmablas_slat2d( uplo[iuplo], n, dSA, ldda, dA, ldda, &info );
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gbytes / gpu_time;
if (info != 0)
printf("magmablas_slat2d returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix( n, n, dA, ldda, R, lda );
if ( opts.verbose ) {
printf( "A= " ); magma_dprint( n, n, A, lda );
printf( "SA= " ); magma_sprint( n, n, SA, lda );
printf( "dA= " ); magma_dprint_gpu( n, n, dA, ldda );
printf( "dSA=" ); magma_sprint_gpu( n, n, dSA, ldda );
}
/* =====================================================================
compute error |A_magma - A_lapack|
should be zero if both are IEEE compliant
=================================================================== */
blasf77_daxpy( &size, &c_neg_one, A, &ione, R, &ione );
error = lapackf77_dlange( "Fro", &n, &n, R, &lda, work );
printf( "slat2d %5s %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
lapack_uplo_const(uplo[iuplo]), (int) n,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
error, (error == 0 ? "ok" : "failed") );
status += ! (error == 0);
TESTING_FREE_CPU( SA );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( SR );
TESTING_FREE_CPU( R );
TESTING_FREE_DEV( dSA );
TESTING_FREE_DEV( dA );
printf( "\n" );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dtrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double magma_error, cublas_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *piv;
magma_err_t err;
double *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2, *LU, *LUT;
double *d_A, *d_B;
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"side = %c, uplo = %c, transA = %c, diag = %c \n", opts.side, opts.uplo, opts.transA, opts.diag );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("==================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
gflops = FLOPS_DTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC( h_A, double, lda*Ak );
TESTING_MALLOC( LU, double, lda*Ak );
TESTING_MALLOC( LUT, double, lda*Ak );
TESTING_MALLOC( h_B, double, ldb*N );
TESTING_MALLOC( h_B1, double, ldb*N );
TESTING_MALLOC( h_X1, double, ldb*N );
TESTING_MALLOC( h_X2, double, ldb*N );
TESTING_MALLOC( h_Bcublas, double, ldb*N );
TESTING_MALLOC( h_Bmagma, double, ldb*N );
TESTING_DEVALLOC( d_A, double, ldda*Ak );
TESTING_DEVALLOC( d_B, double, lddb*N );
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &sizeA, LU );
err = magma_malloc_cpu( (void**) &piv, Ak*sizeof(magma_int_t) ); assert( err == 0 );
lapackf77_dgetrf( &Ak, &Ak, LU, &lda, piv, &info );
int i, j;
for(i=0;i<Ak;i++){
for(j=0;j<Ak;j++){
LUT[j+i*lda] = LU[i+j*lda];
}
}
lapackf77_dlacpy(MagmaUpperStr, &Ak, &Ak, LUT, &lda, LU, &lda);
if(opts.uplo == MagmaLower){
lapackf77_dlacpy(MagmaLowerStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}else{
lapackf77_dlacpy(MagmaUpperStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}
lapackf77_dlarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(double));
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_dsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_dtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_dgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasDtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_dgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_dtrsm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(double));
double alpha2 = MAGMA_D_DIV( c_one, alpha );
blasf77_dtrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_daxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
double norm1 = lapackf77_dlange( "M", &M, &N, h_X1, &ldb, work );
double normx = lapackf77_dlange( "M", &M, &N, h_Bmagma, &ldb, work );
double normA = lapackf77_dlange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(double));
blasf77_dtrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_daxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_dlange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_dlange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_dlange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error );
}
TESTING_FREE( h_A );
TESTING_FREE( LU );
TESTING_FREE( LUT );
TESTING_FREE( h_B );
TESTING_FREE( h_Bcublas );
TESTING_FREE( h_Bmagma );
TESTING_FREE( h_B1 );
TESTING_FREE( h_X1 );
TESTING_FREE( h_X2 );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dormql
*/
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;
magma_int_t ione = 1;
magma_int_t mm, m, n, k, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
double *C, *R, *A, *W, *tau;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf(" M N K side trans CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||QC||_F\n");
printf("===============================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
nb = magma_get_dgeqlf_nb( m );
ldc = m;
// A is m x k (left) or n x k (right)
mm = (side[iside] == MagmaLeft ? m : n);
lda = mm;
gflops = FLOPS_DORMQL( m, n, k, side[iside] ) / 1e9;
if ( side[iside] == MagmaLeft && m < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=left and m < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
if ( side[iside] == MagmaRight && n < k ) {
printf( "%5d %5d %5d %4c %5c skipping because side=right and n < k\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ) );
continue;
}
// need at least 2*nb*nb for geqlf
lwork_max = max( max( m*nb, n*nb ), 2*nb*nb );
TESTING_MALLOC_CPU( C, double, ldc*n );
TESTING_MALLOC_CPU( R, double, ldc*n );
TESTING_MALLOC_CPU( A, double, lda*k );
TESTING_MALLOC_CPU( W, double, lwork_max );
TESTING_MALLOC_CPU( tau, double, k );
// C is full, m x n
size = ldc*n;
lapackf77_dlarnv( &ione, ISEED, &size, C );
lapackf77_dlacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*k;
lapackf77_dlarnv( &ione, ISEED, &size, A );
// compute QL factorization to get Householder vectors in A, tau
magma_dgeqlf( mm, k, A, lda, tau, W, lwork_max, &info );
if (info != 0)
printf("magma_dgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dormql( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, W, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dormql returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_dormql( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
if (info != 0)
printf("magma_dormql (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_D_REAL( W[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("optimal lwork %d > lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
gpu_time = magma_wtime();
magma_dormql( side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, W, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dormql returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_dlange( "Fro", &m, &n, C, &ldc, work );
size = ldc*n;
blasf77_daxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_dlange( "Fro", &m, &n, R, &ldc, work ) / error;
printf( "%5d %5d %5d %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_side_const( side[iside] ),
lapacke_trans_const( trans[itran] ),
cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( W );
TESTING_FREE_CPU( tau );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dpotf2_gpu
*/
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;
magma_int_t N, n2, lda, ldda, info;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_DPOTRF( N ) / 1e9;
if ( N > 512 ) {
printf( "%5d skipping because dpotf2 does not support N > 512\n", (int) N );
continue;
}
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 );
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dpotf2_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dpotf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dpotrf( 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_dpotrf 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*1000., gpu_perf, gpu_time*1000.,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
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;
magma_int_t N, n2, lda, info;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("ngpu %d, uplo %c\n", (int) opts.ngpu, opts.uplo );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
gflops = FLOPS_DPOTRF( N ) / 1e9;
TESTING_MALLOC( h_A, double, n2 );
TESTING_HOSTALLOC( h_R, double, n2 );
/* 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
=================================================================== */
gpu_time = magma_wtime();
magma_dpotrf( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dpotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
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( h_A );
TESTING_HOSTFREE( h_R );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dpotrf
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
double flops, gpu_perf, cpu_perf;
double *h_A, *h_R;
double *d_A;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info;
const char *uplo = MagmaUpperStr;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], matnorm;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
}
if (N>0) size[0] = size[9] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_dpotri_gpu -N %d\n\n", 1024);
}
/* Allocate host memory for the matrix */
n2 = size[9] * size[9];
ldda = ((size[9]+31)/32) * 32;
TESTING_MALLOC( h_A, double, n2);
TESTING_HOSTALLOC( h_R, double, n2);
TESTING_DEVALLOC( d_A, double, ldda*size[9] );
printf(" N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("========================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
flops = FLOPS_DPOTRI( (double)N ) / 1000000;
ldda = ((N+31)/32)*32;
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_D_SET2REAL( h_A[i*lda+i], ( MAGMA_D_REAL(h_A[i*lda+i]) + 1.*N ) );
for(j=0; j<i; j++)
h_A[i*lda+j] = (h_A[j*lda+i]);
}
}
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
//cublasSetMatrix( N, N, sizeof(double), h_A, lda, d_A, ldda);
//magma_dpotrf_gpu(uplo[0], N, d_A, ldda, &info);
/* factorize matrix */
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_dpotrf_gpu(uplo[0], 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 );
start = get_current_time();
magma_dpotri_gpu(uplo[0], N, d_A, ldda, &info);
end = get_current_time();
if (info != 0)
printf("magma_dpotri_gpu returned error %d\n", (int) info);
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_dpotrf(uplo, &N, h_A, &lda, &info);
start = get_current_time();
lapackf77_dpotri(uplo, &N, h_A, &lda, &info);
end = get_current_time();
if (info != 0)
printf("lapackf77_dpotri returned error %d\n", (int) info);
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_dgetmatrix( N, N, d_A, ldda, h_R, lda );
matnorm = lapackf77_dlange("f", &N, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %6.2f %6.2f %e\n",
(int) size[i], cpu_perf, gpu_perf,
lapackf77_dlange("f", &N, &N, h_R, &lda, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dtrmm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
double cublas_error, Cnorm, work[1];
magma_int_t M, N;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double *h_A, *h_B, *h_Bcublas;
magmaDouble_ptr d_A, d_B, d_C;
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% If running lapack (option --lapack), CUBLAS error is computed\n"
"%% relative to CPU BLAS result.\n\n");
printf("%% side = %s, uplo = %s, transA = %s, diag = %s \n",
lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) );
printf("%% M N CUBLAS Gflop/s (ms) CPU Gflop/s (ms) CUBLAS error\n");
printf("%%=================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
gflops = FLOPS_DTRMM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
lddc = lddb;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC_CPU( h_A, double, lda*Ak );
TESTING_MALLOC_CPU( h_B, double, ldb*N );
TESTING_MALLOC_CPU( h_Bcublas, double, ldb*N );
TESTING_MALLOC_DEV( d_A, double, ldda*Ak );
TESTING_MALLOC_DEV( d_B, double, lddb*N );
TESTING_MALLOC_DEV( d_C, double, lddc*N );
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_dlarnv( &ione, ISEED, &sizeB, h_B );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_dsetmatrix( Ak, Ak, h_A, lda, d_A, ldda, opts.queue );
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb, opts.queue );
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
cublas_time = magma_sync_wtime( opts.queue );
#ifdef HAVE_CUBLAS
cublasDtrmm( opts.handle, cublas_side_const(opts.side), cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
M, N,
&alpha, d_A, ldda,
d_B, lddb,
d_C, lddc ); // output C; differs from BLAS standard
#else
magma_dtrmm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, 0, ldda,
d_B, 0, lddb, opts.queue );
#endif
cublas_time = magma_sync_wtime( opts.queue ) - cublas_time;
cublas_perf = gflops / cublas_time;
#ifdef HAVE_CUBLAS
magma_dgetmatrix( M, N, d_C, lddc, h_Bcublas, ldb, opts.queue );
#else
magma_dgetmatrix( M, N, d_B, 0, lddb, h_Bcublas, ldb, opts.queue, opts.queue );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_dtrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_dlange( "M", &M, &N, h_B, &ldb, work );
blasf77_daxpy( &sizeB, &c_neg_one, h_B, &ione, h_Bcublas, &ione );
cublas_error = lapackf77_dlange( "M", &M, &N, h_Bcublas, &ldb, work ) / Cnorm;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) --- ( --- ) --- ---\n",
(int) M, (int) N,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_Bcublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
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_R, *tau, *dtau, *h_work, tmp[1];
double *d_A;
double *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_dgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC( tau, double, min_mn );
TESTING_MALLOC( h_A, double, n2 );
TESTING_HOSTALLOC( h_R, double, n2 );
TESTING_DEVALLOC( d_A, double, ldda*N );
TESTING_DEVALLOC( dtau, double, min_mn );
TESTING_DEVALLOC(dwork, double, min_mn );
TESTING_MALLOC( 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 );
// warmup
magma_dgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_dgeqr2_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_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dgeqrf(&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_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 );
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, error );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( dtau );
TESTING_DEVFREE( dwork );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
int main( int argc, char** argv)
{
real_Double_t gflops, magma_perf, magma_time, clblas_perf, clblas_time, cpu_perf, cpu_time;
double magma_error, clblas_error, work[1];
int transA = MagmaNoTrans;
int transB = MagmaNoTrans;
magma_int_t istart = 1024;
magma_int_t iend = 6240;
magma_int_t M, M0 = 0;
magma_int_t N, N0 = 0;
magma_int_t K, K0 = 0;
magma_int_t i;
magma_int_t Am, An, Bm, Bn;
magma_int_t szeA, szeB, szeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double *h_A, *h_B, *h_C, *h_C2, *h_C3;
magmaDouble_ptr d_A, d_B, d_C;
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
double beta = MAGMA_D_MAKE( -0.48, 0.38 );
int lapack = getenv("MAGMA_RUN_LAPACK") != NULL;
int count = 1;
printf("\nUsage: testing_dgemm [-NN|NT|TN|TT|NC|CN|TC|CT|CC] -M m -N n -K k -count c -l\n"
" -l or setting $MAGMA_RUN_LAPACK runs CPU BLAS,\n"
" and computes both MAGMA and CLBLAS error using CPU BLAS result.\n"
" Else, MAGMA error is computed using CLBLAS result.\n\n");
for( int i = 1; i < argc; ++i ) {
if ( strcmp("-N", argv[i]) == 0 && i+1 < argc ){
N0 = atoi(argv[++i]);
}
else if ( strcmp("-M", argv[i]) == 0 && i+1 < argc ){
M0 = atoi(argv[++i]);
}
else if ( strcmp("-K", argv[i]) == 0 && i+1 < argc ){
K0 = atoi(argv[++i]);
}
else if (strcmp("-NN", argv[i])==0){
transA = transB = MagmaNoTrans;
}
else if (strcmp("-TT", argv[i])==0){
transA = transB = MagmaTrans;
}
else if (strcmp("-NT", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaTrans;
}
else if (strcmp("-TN", argv[i])==0){
transA = MagmaTrans;
transB = MagmaNoTrans;
}
else if (strcmp("-NC", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaTrans;
}
else if (strcmp("-TC", argv[i])==0){
transA = MagmaTrans;
transB = MagmaTrans;
}
else if (strcmp("-CN", argv[i])==0){
transA = MagmaTrans;
transB = MagmaNoTrans;
}
else if (strcmp("-CT", argv[i])==0){
transA = MagmaTrans;
transB = MagmaTrans;
}
else if (strcmp("-CC", argv[i])==0){
transA = transB = MagmaTrans;
}
else if (strcmp("-l", argv[i])==0) {
lapack = true;
}
else if ( strcmp("-count", argv[i]) == 0 && i+1 < argc ){
count = atoi(argv[++i]);
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( (M0 != 0) && (N0 != 0) && (K0 != 0) )
iend = istart + 1;
M = N = K = iend;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if ( K0 != 0 ) K = K0;
if( transA == MagmaNoTrans ) {
Am = M;
An = K;
} else {
Am = K;
An = M;
}
if( transB == MagmaNoTrans ) {
Bm = K;
Bn = N;
} else {
Bm = N;
Bn = K;
}
/* 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);
}
lda = ldc = M;
ldb = Bm;
ldda = ((M+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ldda;
K += 32;
M += 32;
N += 32;
TESTING_MALLOC_CPU( h_A, double, lda*K );
TESTING_MALLOC_CPU( h_B, double, ldb*Bn );
TESTING_MALLOC_CPU( h_C, double, ldc*N );
TESTING_MALLOC_CPU( h_C2, double, ldc*N );
TESTING_MALLOC_CPU( h_C3, double, ldc*N );
TESTING_MALLOC_DEV( d_A, double, ldda*K );
TESTING_MALLOC_DEV( d_B, double, lddb*Bn );
TESTING_MALLOC_DEV( d_C, double, lddc*N );
printf("Testing transA = %c transB = %c\n", *lapack_const(transA), *lapack_const(transB));
printf(" M N K MAGMA Gflop/s (sec) CLBLAS Gflop/s (sec) CPU Gflop/s (sec) MAGMA error CLBLAS error\n");
printf("===========================================================================================================\n");
for( i=istart; i<iend; i = (int)(i*1.25) ) {
for( int cnt = 0; cnt < count; ++cnt ) {
M = N = K = i;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if ( K0 != 0 ) K = K0;
if( transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if( transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
gflops = FLOPS_DGEMM( M, N, K ) / 1e9;
ldc = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
szeA = lda * An;
szeB = ldb * Bn;
szeC = ldc * N;
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &szeA, h_A );
lapackf77_dlarnv( &ione, ISEED, &szeB, h_B );
lapackf77_dlarnv( &ione, ISEED, &szeC, h_C );
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_dsetmatrix( Am, An, h_A, 0, lda, d_A, 0, ldda, queue );
magma_dsetmatrix( Bm, Bn, h_B, 0, ldb, d_B, 0, lddb, queue );
magma_dsetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magmablas_dgemm_reduce( M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_dsetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_queue_sync(queue);
magma_time = magma_wtime();
magmablas_dgemm_reduce( M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_queue_sync(queue);
magma_time = magma_wtime() - magma_time;
magma_perf = gflops / magma_time;
magma_dgetmatrix( M, N, d_C, 0, lddc, h_C2, 0, ldc, queue );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_dsetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_dgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_dsetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_queue_sync(queue);
clblas_time = magma_wtime();
magma_dgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_queue_sync(queue);
clblas_time = magma_wtime() - clblas_time;
clblas_perf = gflops / clblas_time;
magma_dgetmatrix( M, N, d_C, 0, lddc, h_C3, 0, ldc, queue );
/* =====================================================================
Performs operation using BLAS
=================================================================== */
if ( lapack ) {
cpu_time = magma_wtime();
blasf77_dgemm( lapack_const(transA), lapack_const(transB), &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Error Computation and Performance Compariosn
=================================================================== */
if ( lapack ) {
// compare both magma & clblas to lapack
blasf77_daxpy(&szeC, &c_neg_one, h_C, &ione, h_C2, &ione);
magma_error = lapackf77_dlange("M", &M, &N, h_C2, &ldc, work);
blasf77_daxpy(&szeC, &c_neg_one, h_C, &ione, h_C3, &ione);
clblas_error = lapackf77_dlange("M", &M, &N, h_C3, &ldc, work);
printf("%5d %5d %5d %7.2f (%7.4f) %7.2f (%7.4f) %7.2f (%7.4f) %8.2e %8.2e\n",
(int) M, (int) N, (int) K,
magma_perf, magma_time, clblas_perf, clblas_time, cpu_perf, cpu_time,
magma_error, clblas_error );
}
else {
// compare magma to clblas
blasf77_daxpy(&szeC, &c_neg_one, h_C3, &ione, h_C2, &ione);
magma_error = lapackf77_dlange("M", &M, &N, h_C2, &ldc, work);
printf("%5d %5d %5d %7.2f (%7.4f) %7.2f (%7.4f) --- ( --- ) %8.2e ---\n",
(int) M, (int) N, (int) K,
magma_perf, magma_time, clblas_perf, clblas_time,
magma_error );
}
}
if ( count > 1 ) {
printf( "\n" );
}
}
/* Memory clean up */
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_C2 );
TESTING_FREE_CPU( h_C3 );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
magma_queue_destroy( queue );
magma_finalize();
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
double *h_A, *h_R;
magmaDouble_ptr d_lA[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, info;
double mz_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], matnorm, diffnorm;
magma_int_t num_gpus0 = 1, num_gpus, flag = 0;
int nb, mb, 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;
}else exit(1);
}
if(strcmp("-NGPU", argv[i])==0)
num_gpus0 = atoi(argv[++i]);
if(strcmp("-UPLO", argv[i])==0){
if(strcmp("L", argv[++i])==0){
uplo = MagmaLower;
}else{
uplo = MagmaUpper;
}
}
}
}
else {
printf("\nUsage: \n");
printf(" testing_dpotrf_mgpu -N %d -NGPU %d -UPLO -L\n\n", 1024, num_gpus0);
}
/* looking for max. ldda */
ldda = 0;
n2 = 0;
for(i=0;i<10;i++){
N = size[i];
nb = magma_get_dpotrf_nb(N);
mb = nb;
if(num_gpus0 > N/nb){
num_gpus = N/nb;
if(N%nb != 0) num_gpus ++;
}else{
num_gpus = num_gpus0;
}
n_local = nb*(1+N/(nb*num_gpus))*mb*((N+mb-1)/mb);
if(n_local > ldda) ldda = n_local;
if(n2 < N*N) n2 = N*N;
if(flag != 0) break;
}
/* Allocate host memory for the matrix */
TESTING_MALLOC_PIN( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
/* Initialize */
magma_queue_t queues[MagmaMaxGPUs * 2];
//magma_queue_t queues[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_gpus;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("each buffer size: %d\n", ldda);
/* allocate local matrix on Buffers */
for(i=0; i<num_gpus0; i++){
TESTING_MALLOC_DEV( d_lA[i], double, ldda );
}
printf("\n\n");
printf("Using GPUs: %d\n", num_gpus0);
if(uplo == MagmaUpper){
printf("\n testing_dpotrf_mgpu -N %d -NGPU %d -UPLO U\n\n", N, num_gpus0);
}else{
printf("\n testing_dpotrf_mgpu -N %d -NGPU %d -UPLO L\n\n", N, num_gpus0);
}
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma-R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS( (double)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
for( int i = 0; i < N; ++i ) {
h_A(i,i) = MAGMA_D_MAKE( MAGMA_D_REAL(h_A(i,i)) + N, 0 );
for( int j = 0; j < i; ++j ) {
h_A(i, j) = MAGMA_D_CNJG( h_A(j,i) );
}
}
lapackf77_dlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
/* Warm up to measure the performance */
nb = magma_get_dpotrf_nb(N);
if(num_gpus0 > N/nb){
num_gpus = N/nb;
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;
}
/* distribute matrix to gpus */
if(uplo == MagmaUpper){
// Upper
ldda = ((N+mb-1)/mb)*mb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_dsetmatrix( N, nk,
&h_A[j*lda], lda,
d_lA[k], j/(nb*num_gpus)*nb*ldda, ldda,
queues[2*k]);
}
}else{
// Lower
ldda = (1+N/(nb*num_gpus))*nb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_dsetmatrix( nk, N, &h_A[j], lda,
d_lA[k], (j/(nb*num_gpus)*nb), ldda,
queues[2*k]);
}
}
magma_dpotrf_mgpu( num_gpus, uplo, N, d_lA, 0, ldda, queues, &info );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* distribute matrix to gpus */
if(uplo == MagmaUpper){
// Upper
ldda = ((N+mb-1)/mb)*mb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_dsetmatrix( N, nk,
&h_A[j*lda], lda,
d_lA[k], j/(nb*num_gpus)*nb*ldda, ldda,
queues[2*k]);
}
}else{
// Lower
ldda = (1+N/(nb*num_gpus))*nb;
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_dsetmatrix( nk, N, &h_A[j], lda,
d_lA[k], (j/(nb*num_gpus)*nb), ldda,
queues[2*k]);
}
}
gpu_time = magma_wtime();
magma_dpotrf_mgpu( num_gpus, uplo, N, d_lA, 0, ldda, queues, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf( "magma_dpotrf had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* gather matrix from gpus */
if(uplo==MagmaUpper){
// Upper
for(j=0;j<N;j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_dgetmatrix( N, nk,
d_lA[k], j/(nb*num_gpus)*nb*ldda, ldda,
&h_R[j*lda], lda, queues[2*k]);
}
}else{
// Lower
for(j=0; j<N; j+=nb){
k = (j/nb)%num_gpus;
nk = min(nb, N-j);
magma_dgetmatrix( nk, N,
d_lA[k], (j/(nb*num_gpus)*nb), ldda,
&h_R[j], lda, queues[2*k] );
}
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
if(uplo == MagmaLower){
lapackf77_dpotrf( MagmaLowerStr, &N, h_A, &lda, &info );
}else{
lapackf77_dpotrf( MagmaUpperStr, &N, h_A, &lda, &info );
}
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf( "lapackf77_dpotrf had error %d.\n", info );
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result compared to LAPACK
|R_magma - R_lapack| / |R_lapack|
=================================================================== */
matnorm = lapackf77_dlange("f", &N, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &mz_one, h_A, &ione, h_R, &ione);
diffnorm = lapackf77_dlange("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 );
if (flag != 0)
break;
}
/* clean up */
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_R );
for(i=0;i<num_gpus;i++){
TESTING_FREE_DEV( d_lA[i] );
magma_queue_destroy( queues[2*i] );
magma_queue_destroy( queues[2*i+1] );
}
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgesv
*/
int main(int argc , char **argv)
{
TESTING_CUDA_INIT();
real_Double_t gflops, gpu_perf, gpu_time;
double Rnorm, Anorm, Xnorm, *work;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_LU, *h_B, *h_X;
magma_int_t *ipiv;
magma_int_t lda, ldb;
magma_int_t i, info, szeA, szeB;
magma_int_t ione = 1;
magma_int_t N = 0;
magma_int_t NRHS = 100;
magma_int_t ISEED[4] = {0,0,0,1};
const int MAXTESTS = 10;
magma_int_t size[MAXTESTS] = { 1024, 2048, 3072, 4032, 5184, 6016, 7040, 8064, 9088, 10112 };
// process command line arguments
printf( "\nUsage: %s -N <matrix size> -R <right hand sides>\n", argv[0] );
printf( " -N can be repeated up to %d times\n\n", MAXTESTS );
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 );
size[ntest] = atoi( argv[++i] );
magma_assert( size[ntest] > 0, "error: -N %s is invalid; must be > 0.\n", argv[i] );
N = max( N, size[ntest] );
ntest++;
}
else if ( strcmp("-R", argv[i]) == 0 && i+1 < argc ) {
NRHS = atoi( argv[++i] );
magma_assert( NRHS > 0, "error: -R %is is invalid; must be > 0.\n", argv[i] );
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( ntest == 0 ) {
ntest = MAXTESTS;
N = size[ntest-1];
}
// allocate maximum amount of memory required
lda = ldb = N;
TESTING_MALLOC( h_A, double, lda*N );
TESTING_MALLOC( h_LU, double, lda*N );
TESTING_MALLOC( h_B, double, ldb*NRHS );
TESTING_MALLOC( h_X, double, ldb*NRHS );
TESTING_MALLOC( work, double, N );
TESTING_MALLOC( ipiv, magma_int_t, N );
printf(" N NRHS GPU GFlop/s (sec) ||B - AX|| / ||A||*||X||\n");
printf("===========================================================\n");
for( i = 0; i < ntest; ++i ) {
N = size[i];
lda = ldb = N;
gflops = ( FLOPS_DGETRF( (double)N, (double)N ) +
FLOPS_DGETRS( (double)N, (double)NRHS ) ) / 1e9;
/* Initialize the matrices */
szeA = lda*N;
szeB = ldb*NRHS;
lapackf77_dlarnv( &ione, ISEED, &szeA, h_A );
lapackf77_dlarnv( &ione, ISEED, &szeB, h_B );
// copy A to LU and B to X; save A and B for residual
lapackf77_dlacpy( "F", &N, &N, h_A, &lda, h_LU, &lda );
lapackf77_dlacpy( "F", &N, &NRHS, h_B, &ldb, h_X, &ldb );
//=====================================================================
// Solve Ax = b through an LU factorization
//=====================================================================
gpu_time = magma_wtime();
magma_dgesv( N, NRHS, h_LU, lda, ipiv, h_X, ldb, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_dgesv returned error %d.\n", (int) info);
gpu_perf = gflops / gpu_time;
//=====================================================================
// Residual
//=====================================================================
Anorm = lapackf77_dlange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_dlange("I", &N, &NRHS, h_X, &ldb, work);
blasf77_dgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &NRHS, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb);
Rnorm = lapackf77_dlange("I", &N, &NRHS, h_B, &ldb, work);
printf( "%5d %5d %7.2f (%7.2f) %8.2e\n",
(int) N, (int) NRHS, gpu_perf, gpu_time, Rnorm/(Anorm*Xnorm) );
}
/* Memory clean up */
TESTING_FREE( h_A );
TESTING_FREE( h_LU );
TESTING_FREE( h_B );
TESTING_FREE( h_X );
TESTING_FREE( work );
TESTING_FREE( ipiv );
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgesdd (SVD with Divide & Conquer)
Please keep code in testing_zgesdd.cpp and testing_zgesvd.cpp similar.
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
magmaDoubleComplex *h_A, *h_R, *U, *VT, *h_work;
magmaDoubleComplex dummy[1];
double *S1, *S2;
#ifdef COMPLEX
magma_int_t lrwork=0;
double *rwork;
#endif
magma_int_t *iwork;
magma_int_t M, N, N_U, M_VT, lda, ldu, ldv, n2, min_mn, max_mn, info, nb, lwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_vec_t jobz;
magma_int_t status = 0;
MAGMA_UNUSED( max_mn ); // used only in complex
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
jobz = opts.jobu;
magma_vec_t jobs[] = { MagmaNoVec, MagmaSomeVec, MagmaOverwriteVec, MagmaAllVec };
if ( opts.check && ! opts.all && (jobz == MagmaNoVec)) {
printf( "%% NOTE: some checks require that singular vectors are computed;\n"
"%% set jobz (option -U[NASO]) to be S, O, or A.\n\n" );
}
printf("%% jobz M N CPU time (sec) GPU time (sec) |S1-S2| |A-USV^H| |I-UU^H|/M |I-VV^H|/N S sorted\n");
printf("%%==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int ijobz = 0; ijobz < 4; ++ijobz ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
if ( opts.all ) {
jobz = jobs[ ijobz ];
}
else if ( ijobz > 0 ) {
// if not testing all, run only once, when ijobz = 0,
// but jobz come from opts (above loops).
continue;
}
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
max_mn = max(M, N);
N_U = (jobz == MagmaAllVec ? M : min_mn);
M_VT = (jobz == MagmaAllVec ? N : min_mn);
lda = M;
ldu = M;
ldv = M_VT;
n2 = lda*N;
nb = magma_get_zgesvd_nb( M, N );
// x and y abbreviations used in zgesdd and dgesdd documentation
magma_int_t x = max(M,N);
magma_int_t y = min(M,N);
#ifdef COMPLEX
bool tall = (x >= int(y*17/9.)); // true if tall (m >> n) or wide (n >> m)
#else
bool tall = (x >= int(y*11/6.)); // true if tall (m >> n) or wide (n >> m)
#endif
// query or use formula for workspace size
switch( opts.svd_work ) {
case 0: { // query for workspace size
lwork = -1;
magma_zgesdd( jobz, M, N,
NULL, lda, NULL, NULL, ldu, NULL, ldv, dummy, lwork,
#ifdef COMPLEX
NULL,
#endif
NULL, &info );
lwork = (int) MAGMA_Z_REAL( dummy[0] );
break;
}
case 1: // minimum
case 2: // optimal
case 3: { // optimal (for gesdd, 2 & 3 are same; for gesvd, they differ)
// formulas from zgesdd and dgesdd documentation
bool sml = (opts.svd_work == 1); // 1 is small workspace, 2,3 are large workspace
#ifdef COMPLEX // ----------------------------------------
if (jobz == MagmaNoVec) {
if (tall) { lwork = 2*y + (2*y)*nb; }
else { lwork = 2*y + (x+y)*nb; }
}
if (jobz == MagmaOverwriteVec) {
if (tall) {
if (sml) { lwork = 2*y*y + 2*y + (2*y)*nb; }
else { lwork = y*y + x*y + 2*y + (2*y)*nb; } // not big deal
}
else {
//if (sml) { lwork = 2*y + max( (x+y)*nb, y*y + y ); }
//else { lwork = 2*y + max( (x+y)*nb, x*y + y*nb ); }
// LAPACK 3.4.2 over-estimates workspaces. For compatability, use these:
if (sml) { lwork = 2*y + max( (x+y)*nb, y*y + x ); }
else { lwork = 2*y + (x+y)*nb + x*y; }
}
}
if (jobz == MagmaSomeVec) {
if (tall) { lwork = y*y + 2*y + (2*y)*nb; }
else { lwork = 2*y + (x+y)*nb; }
}
if (jobz == MagmaAllVec) {
if (tall) {
if (sml) { lwork = y*y + 2*y + max( (2*y)*nb, x ); }
else { lwork = y*y + 2*y + max( (2*y)*nb, x*nb ); }
}
else { lwork = 2*y + (x+y)*nb; }
}
#else // REAL ----------------------------------------
if (jobz == MagmaNoVec) {
if (tall) { lwork = 3*y + max( (2*y)*nb, 7*y ); }
else { lwork = 3*y + max( (x+y)*nb, 7*y ); }
}
if (jobz == MagmaOverwriteVec) {
if (tall) {
if (sml) { lwork = y*y + 3*y + max( (2*y)*nb, 4*y*y + 4*y ); }
else { lwork = y*y + 3*y + max( max( (2*y)*nb, 4*y*y + 4*y ), y*y + y*nb ); }
}
else {
if (sml) { lwork = 3*y + max( (x+y)*nb, 4*y*y + 4*y ); }
else { lwork = 3*y + max( (x+y)*nb, 3*y*y + 4*y + x*y ); } // extra space not too important?
}
}
if (jobz == MagmaSomeVec) {
if (tall) { lwork = y*y + 3*y + max( (2*y)*nb, 3*y*y + 4*y ); }
else { lwork = 3*y + max( (x+y)*nb, 3*y*y + 4*y ); }
}
if (jobz == MagmaAllVec) {
if (tall) {
if (sml) { lwork = y*y + max( 3*y + max( (2*y)*nb, 3*y*y + 4*y ), y + x ); }
else { lwork = y*y + max( 3*y + max( (2*y)*nb, 3*y*y + 4*y ), y + x*nb ); }
// LAPACK 3.4.2 over-estimates workspaces. For compatability, use these:
//if (sml) { lwork = y*y + 3*y + max( (2*y)*nb, 3*y*y + 3*y + x ); }
//else { lwork = y*y + max( 3*y + max( (2*y)*nb, max( 3*y*y + 3*y + x, 3*y*y + 4*y )), y + x*nb ); }
}
else { lwork = 3*y + max( (x+y)*nb, 3*y*y + 4*y ); }
}
#endif
break;
}
default: {
fprintf( stderr, "Error: unknown option svd_work %d\n", (int) opts.svd_work );
return -1;
break;
}
}
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( VT, magmaDoubleComplex, ldv*N ); // N x N (jobz=A) or min(M,N) x N
TESTING_MALLOC_CPU( U, magmaDoubleComplex, ldu*N_U ); // M x M (jobz=A) or M x min(M,N)
TESTING_MALLOC_CPU( S1, double, min_mn );
TESTING_MALLOC_CPU( S2, double, min_mn );
TESTING_MALLOC_CPU( iwork, magma_int_t, 8*min_mn );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
#ifdef COMPLEX
if (jobz == MagmaNoVec) {
// requires 5*min_mn, but MKL (11.1) seems to have a bug
// requiring 7*min_mn in some cases (e.g., jobz=N, m=100, n=170)
lrwork = 7*min_mn;
}
else if (tall) {
lrwork = 5*min_mn*min_mn + 5*min_mn;
}
else {
lrwork = max( 5*min_mn*min_mn + 5*min_mn,
2*max_mn*min_mn + 2*min_mn*min_mn + min_mn );
}
TESTING_MALLOC_CPU( rwork, double, lrwork );
#endif
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgesdd( jobz, M, N,
h_R, lda, S1, U, ldu, VT, ldv, h_work, lwork,
#ifdef COMPLEX
rwork,
#endif
iwork, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0) {
printf("magma_zgesdd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
double eps = lapackf77_dlamch( "E" );
double result[5] = { -1/eps, -1/eps, -1/eps, -1/eps, -1/eps };
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zcds]drvbd routine.
A is factored as A = U diag(S) VT and the following 4 tests computed:
(1) | A - U diag(S) VT | / ( |A| max(M,N) )
(2) | I - U^H U | / ( M )
(3) | I - VT VT^H | / ( N )
(4) S contains MNMIN nonnegative values in decreasing order.
(Return 0 if true, 1/ULP if false.)
=================================================================== */
magma_int_t izero = 0;
// get size and location of U and V^T depending on jobz
// U2=NULL and VT2=NULL if they were not computed (e.g., jobz=N)
magmaDoubleComplex *U2 = NULL;
magmaDoubleComplex *VT2 = NULL;
if ( jobz == MagmaSomeVec || jobz == MagmaAllVec ) {
U2 = U;
VT2 = VT;
}
else if ( jobz == MagmaOverwriteVec ) {
if ( M >= N ) {
U2 = h_R;
ldu = lda;
VT2 = VT;
}
else {
U2 = U;
VT2 = h_R;
ldv = lda;
}
}
// zbdt01 needs M+N
// zunt01 prefers N*(N+1) to check U; M*(M+1) to check V
magma_int_t lwork_err = M+N;
if ( U2 != NULL ) {
lwork_err = max( lwork_err, N_U*(N_U+1) );
}
if ( VT2 != NULL ) {
lwork_err = max( lwork_err, M_VT*(M_VT+1) );
}
magmaDoubleComplex *h_work_err;
TESTING_MALLOC_CPU( h_work_err, magmaDoubleComplex, lwork_err );
// zbdt01 and zunt01 need max(M,N), depending
double *rwork_err;
TESTING_MALLOC_CPU( rwork_err, double, max(M,N) );
if ( U2 != NULL && VT2 != NULL ) {
// since KD=0 (3rd arg), E is not referenced so pass NULL (9th arg)
lapackf77_zbdt01(&M, &N, &izero, h_A, &lda,
U2, &ldu, S1, NULL, VT2, &ldv,
h_work_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[0]);
}
if ( U2 != NULL ) {
lapackf77_zunt01("Columns", &M, &N_U, U2, &ldu, h_work_err, &lwork_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[1]);
}
if ( VT2 != NULL ) {
lapackf77_zunt01("Rows", &M_VT, &N, VT2, &ldv, h_work_err, &lwork_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[2]);
}
result[3] = 0.;
for (int j=0; j < min_mn-1; j++) {
if ( S1[j] < S1[j+1] )
result[3] = 1.;
if ( S1[j] < 0. )
result[3] = 1.;
}
if (min_mn > 1 && S1[min_mn-1] < 0.)
result[3] = 1.;
result[0] *= eps;
result[1] *= eps;
result[2] *= eps;
TESTING_FREE_CPU( h_work_err );
TESTING_FREE_CPU( rwork_err );
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zgesdd( lapack_vec_const(jobz), &M, &N,
h_A, &lda, S2, U, &ldu, VT, &ldv, h_work, &lwork,
#ifdef COMPLEX
rwork,
#endif
iwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info != 0) {
printf("lapackf77_zgesdd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
double work[1], c_neg_one = -1;
blasf77_daxpy(&min_mn, &c_neg_one, S1, &ione, S2, &ione);
result[4] = lapackf77_dlange("f", &min_mn, &ione, S2, &min_mn, work);
result[4] /= lapackf77_dlange("f", &min_mn, &ione, S1, &min_mn, work);
printf(" %c %5d %5d %7.2f %7.2f %8.2e",
lapack_vec_const(jobz)[0],
(int) M, (int) N, cpu_time, gpu_time, result[4] );
}
else {
printf(" %c %5d %5d --- %7.2f --- ",
lapack_vec_const(jobz)[0],
(int) M, (int) N, gpu_time );
}
if ( opts.check ) {
if ( result[0] < 0. ) { printf(" --- "); } else { printf(" %#9.3g", result[0]); }
if ( result[1] < 0. ) { printf(" --- "); } else { printf(" %#9.3g", result[1]); }
if ( result[2] < 0. ) { printf(" --- "); } else { printf(" %#9.3g", result[2]); }
bool okay = (result[0] < tol) && (result[1] < tol) && (result[2] < tol) && (result[3] == 0.) && (result[4] < tol);
printf(" %3s %s\n", (result[3] == 0. ? "yes" : "no"), (okay ? "ok" : "failed"));
status += ! okay;
}
else {
printf("\n");
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( VT );
TESTING_FREE_CPU( U );
TESTING_FREE_CPU( S1 );
TESTING_FREE_CPU( S2 );
TESTING_FREE_CPU( iwork );
#ifdef COMPLEX
TESTING_FREE_CPU( rwork );
#endif
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
fflush( stdout );
}}
if ( opts.all || opts.niter > 1 ) {
printf("\n");
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
DGEEV computes for an N-by-N real nonsymmetric matrix A, the
eigenvalues and, optionally, the left and/or right eigenvectors.
The right eigenvector v(j) of A satisfies
A * v(j) = lambda(j) * v(j)
where lambda(j) is its eigenvalue.
The left eigenvector u(j) of A satisfies
u(j)**T * A = lambda(j) * u(j)**T
where u(j)**T denotes the transpose of u(j).
The computed eigenvectors are normalized to have Euclidean norm
equal to 1 and largest component real.
Arguments
---------
@param[in]
jobvl magma_vec_t
- = MagmaNoVec: left eigenvectors of A are not computed;
- = MagmaVec: left eigenvectors of are computed.
@param[in]
jobvr magma_vec_t
- = MagmaNoVec: right eigenvectors of A are not computed;
- = MagmaVec: right eigenvectors of A are computed.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N)
On entry, the N-by-N matrix A.
On exit, A has been overwritten.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
wr DOUBLE PRECISION array, dimension (N)
@param[out]
wi DOUBLE PRECISION array, dimension (N)
WR and WI contain the real and imaginary parts,
respectively, of the computed eigenvalues. Complex
conjugate pairs of eigenvalues appear consecutively
with the eigenvalue having the positive imaginary part
first.
@param[out]
VL DOUBLE PRECISION array, dimension (LDVL,N)
If JOBVL = MagmaVec, the left eigenvectors u(j) are stored one
after another in the columns of VL, in the same order
as their eigenvalues.
If JOBVL = MagmaNoVec, VL is not referenced.
u(j) = VL(:,j), the j-th column of VL.
@param[in]
ldvl INTEGER
The leading dimension of the array VL. LDVL >= 1; if
JOBVL = MagmaVec, LDVL >= N.
@param[out]
VR DOUBLE PRECISION array, dimension (LDVR,N)
If JOBVR = MagmaVec, the right eigenvectors v(j) are stored one
after another in the columns of VR, in the same order
as their eigenvalues.
If JOBVR = MagmaNoVec, VR is not referenced.
v(j) = VR(:,j), the j-th column of VR.
@param[in]
ldvr INTEGER
The leading dimension of the array VR. LDVR >= 1; if
JOBVR = MagmaVec, LDVR >= N.
@param[out]
work (workspace) DOUBLE PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= (2+nb)*N.
For optimal performance, LWORK >= (2+2*nb)*N.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = i, the QR algorithm failed to compute all the
eigenvalues, and no eigenvectors have been computed;
elements and i+1:N of W contain eigenvalues which have
converged.
@ingroup magma_dgeev_driver
********************************************************************/
extern "C" magma_int_t
magma_dgeev_m(
magma_vec_t jobvl, magma_vec_t jobvr, magma_int_t n,
double *A, magma_int_t lda,
double *wr, double *wi,
double *VL, magma_int_t ldvl,
double *VR, magma_int_t ldvr,
double *work, magma_int_t lwork,
magma_int_t *info )
{
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
const magma_int_t ione = 1;
const magma_int_t izero = 0;
double d__1, d__2;
double r, cs, sn, scl;
double dum[1], eps;
double anrm, cscale, bignum, smlnum;
magma_int_t i, k, ilo, ihi;
magma_int_t ibal, ierr, itau, iwrk, nout, liwrk, nb;
magma_int_t scalea, minwrk, optwrk, lquery, wantvl, wantvr, select[1];
magma_side_t side = MagmaRight;
magma_timer_t time_total=0, time_gehrd=0, time_unghr=0, time_hseqr=0, time_trevc=0, time_sum=0;
magma_flops_t flop_total=0, flop_gehrd=0, flop_unghr=0, flop_hseqr=0, flop_trevc=0, flop_sum=0;
timer_start( time_total );
flops_start( flop_total );
*info = 0;
lquery = (lwork == -1);
wantvl = (jobvl == MagmaVec);
wantvr = (jobvr == MagmaVec);
if (! wantvl && jobvl != MagmaNoVec) {
*info = -1;
} else if (! wantvr && jobvr != MagmaNoVec) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if ( (ldvl < 1) || (wantvl && (ldvl < n))) {
*info = -9;
} else if ( (ldvr < 1) || (wantvr && (ldvr < n))) {
*info = -11;
}
/* Compute workspace */
nb = magma_get_dgehrd_nb( n );
if (*info == 0) {
minwrk = (2 + nb)*n;
optwrk = (2 + 2*nb)*n;
work[0] = MAGMA_D_MAKE( (double) optwrk, 0. );
if (lwork < minwrk && ! lquery) {
*info = -13;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
#if defined(Version3) || defined(Version4) || defined(Version5)
double *dT;
if (MAGMA_SUCCESS != magma_dmalloc( &dT, nb*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#endif
#if defined(Version4) || defined(Version5)
double *T;
if (MAGMA_SUCCESS != magma_dmalloc_cpu( &T, nb*n )) {
magma_free( dT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
/* Get machine constants */
eps = lapackf77_dlamch( "P" );
smlnum = lapackf77_dlamch( "S" );
bignum = 1. / smlnum;
lapackf77_dlabad( &smlnum, &bignum );
smlnum = magma_dsqrt( smlnum ) / eps;
bignum = 1. / smlnum;
/* Scale A if max element outside range [SMLNUM,BIGNUM] */
anrm = lapackf77_dlange( "M", &n, &n, A, &lda, dum );
scalea = 0;
if (anrm > 0. && anrm < smlnum) {
scalea = 1;
cscale = smlnum;
} else if (anrm > bignum) {
scalea = 1;
cscale = bignum;
}
if (scalea) {
lapackf77_dlascl( "G", &izero, &izero, &anrm, &cscale, &n, &n, A, &lda, &ierr );
}
/* Balance the matrix
* (Workspace: need N)
* - this space is reserved until after gebak */
ibal = 0;
lapackf77_dgebal( "B", &n, A, &lda, &ilo, &ihi, &work[ibal], &ierr );
/* Reduce to upper Hessenberg form
* (Workspace: need 3*N, prefer 2*N + N*NB)
* - including N reserved for gebal/gebak, unused by dgehrd */
itau = ibal + n;
iwrk = itau + n;
liwrk = lwork - iwrk;
timer_start( time_gehrd );
flops_start( flop_gehrd );
#if defined(Version1)
// Version 1 - LAPACK
lapackf77_dgehrd( &n, &ilo, &ihi, A, &lda,
&work[itau], &work[iwrk], &liwrk, &ierr );
#elif defined(Version2)
// Version 2 - LAPACK consistent HRD
magma_dgehrd2( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, &ierr );
#elif defined(Version3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored,
magma_dgehrd( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, dT, &ierr );
#elif defined(Version4) || defined(Version5)
// Version 4 - Multi-GPU, T on host
magma_dgehrd_m( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, T, &ierr );
magma_dsetmatrix( nb, n, T, nb, dT, nb );
#endif
time_sum += timer_stop( time_gehrd );
flop_sum += flops_stop( flop_gehrd );
if (wantvl) {
/* Want left eigenvectors
* Copy Householder vectors to VL */
side = MagmaLeft;
lapackf77_dlacpy( MagmaLowerStr, &n, &n, A, &lda, VL, &ldvl );
/* Generate orthogonal matrix in VL
* (Workspace: need 3*N-1, prefer 2*N + (N-1)*NB)
* - including N reserved for gebal/gebak, unused by dorghr */
timer_start( time_unghr );
flops_start( flop_unghr );
#if defined(Version1) || defined(Version2)
// Version 1 & 2 - LAPACK
lapackf77_dorghr( &n, &ilo, &ihi, VL, &ldvl, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(Version3) || defined(Version4)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_dorghr( n, ilo, ihi, VL, ldvl, &work[itau], dT, nb, &ierr );
#elif defined(Version5)
// Version 5 - Multi-GPU, T on host
magma_dorghr_m( n, ilo, ihi, VL, ldvl, &work[itau], T, nb, &ierr );
#endif
time_sum += timer_stop( time_unghr );
flop_sum += flops_stop( flop_unghr );
timer_start( time_hseqr );
flops_start( flop_hseqr );
/* Perform QR iteration, accumulating Schur vectors in VL
* (Workspace: need N+1, prefer N+HSWORK (see comments) )
* - including N reserved for gebal/gebak, unused by dhseqr */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_dhseqr( "S", "V", &n, &ilo, &ihi, A, &lda, wr, wi,
VL, &ldvl, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
if (wantvr) {
/* Want left and right eigenvectors
* Copy Schur vectors to VR */
side = MagmaBothSides;
lapackf77_dlacpy( "F", &n, &n, VL, &ldvl, VR, &ldvr );
}
}
else if (wantvr) {
/* Want right eigenvectors
* Copy Householder vectors to VR */
side = MagmaRight;
lapackf77_dlacpy( "L", &n, &n, A, &lda, VR, &ldvr );
/* Generate orthogonal matrix in VR
* (Workspace: need 3*N-1, prefer 2*N + (N-1)*NB)
* - including N reserved for gebal/gebak, unused by dorghr */
timer_start( time_unghr );
flops_start( flop_unghr );
#if defined(Version1) || defined(Version2)
// Version 1 & 2 - LAPACK
lapackf77_dorghr( &n, &ilo, &ihi, VR, &ldvr, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(Version3) || defined(Version4)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_dorghr( n, ilo, ihi, VR, ldvr, &work[itau], dT, nb, &ierr );
#elif defined(Version5)
// Version 5 - Multi-GPU, T on host
magma_dorghr_m( n, ilo, ihi, VR, ldvr, &work[itau], T, nb, &ierr );
#endif
time_sum += timer_stop( time_unghr );
flop_sum += flops_stop( flop_unghr );
/* Perform QR iteration, accumulating Schur vectors in VR
* (Workspace: need N+1, prefer N+HSWORK (see comments) )
* - including N reserved for gebal/gebak, unused by dhseqr */
timer_start( time_hseqr );
flops_start( flop_hseqr );
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_dhseqr( "S", "V", &n, &ilo, &ihi, A, &lda, wr, wi,
VR, &ldvr, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
}
else {
/* Compute eigenvalues only
* (Workspace: need N+1, prefer N+HSWORK (see comments) )
* - including N reserved for gebal/gebak, unused by dhseqr */
timer_start( time_hseqr );
flops_start( flop_hseqr );
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_dhseqr( "E", "N", &n, &ilo, &ihi, A, &lda, wr, wi,
VR, &ldvr, &work[iwrk], &liwrk, info );
time_sum += timer_stop( time_hseqr );
flop_sum += flops_stop( flop_hseqr );
}
/* If INFO > 0 from DHSEQR, then quit */
if (*info > 0) {
goto CLEANUP;
}
timer_start( time_trevc );
flops_start( flop_trevc );
if (wantvl || wantvr) {
/* Compute left and/or right eigenvectors
* (Workspace: need 4*N, prefer (2 + 2*nb)*N)
* - including N reserved for gebal/gebak, unused by dtrevc */
liwrk = lwork - iwrk;
#if TREVC_VERSION == 1
lapackf77_dtrevc( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &ierr );
#elif TREVC_VERSION == 2
lapackf77_dtrevc3( lapack_side_const(side), "B", select, &n, A, &lda, VL, &ldvl,
VR, &ldvr, &n, &nout, &work[iwrk], &liwrk, &ierr );
#elif TREVC_VERSION == 3
magma_dtrevc3( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &ierr );
#elif TREVC_VERSION == 4
magma_dtrevc3_mt( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &ierr );
#elif TREVC_VERSION == 5
magma_dtrevc3_mt_gpu( side, MagmaBacktransVec, select, n, A, lda, VL, ldvl,
VR, ldvr, n, &nout, &work[iwrk], liwrk, &ierr );
#else
#error Unknown TREVC_VERSION
#endif
}
time_sum += timer_stop( time_trevc );
flop_sum += flops_stop( flop_trevc );
if (wantvl) {
/* Undo balancing of left eigenvectors
* (Workspace: need N) */
lapackf77_dgebak( "B", "L", &n, &ilo, &ihi, &work[ibal], &n,
VL, &ldvl, &ierr );
/* Normalize left eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
if ( wi[i] == 0. ) {
scl = 1. / magma_cblas_dnrm2( n, VL(0,i), 1 );
blasf77_dscal( &n, &scl, VL(0,i), &ione );
}
else if ( wi[i] > 0. ) {
d__1 = magma_cblas_dnrm2( n, VL(0,i), 1 );
d__2 = magma_cblas_dnrm2( n, VL(0,i+1), 1 );
scl = 1. / lapackf77_dlapy2( &d__1, &d__2 );
blasf77_dscal( &n, &scl, VL(0,i), &ione );
blasf77_dscal( &n, &scl, VL(0,i+1), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = *VL(k,i);
d__2 = *VL(k,i+1);
work[iwrk + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_idamax( &n, &work[iwrk], &ione ) - 1; // subtract 1; k is 0-based
lapackf77_dlartg( VL(k,i), VL(k,i+1), &cs, &sn, &r );
blasf77_drot( &n, VL(0,i), &ione, VL(0,i+1), &ione, &cs, &sn );
*VL(k,i+1) = 0.;
}
}
}
if (wantvr) {
/* Undo balancing of right eigenvectors
* (Workspace: need N) */
lapackf77_dgebak( "B", "R", &n, &ilo, &ihi, &work[ibal], &n,
VR, &ldvr, &ierr );
/* Normalize right eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
if ( wi[i] == 0. ) {
scl = 1. / magma_cblas_dnrm2( n, VR(0,i), 1 );
blasf77_dscal( &n, &scl, VR(0,i), &ione );
}
else if ( wi[i] > 0. ) {
d__1 = magma_cblas_dnrm2( n, VR(0,i), 1 );
d__2 = magma_cblas_dnrm2( n, VR(0,i+1), 1 );
scl = 1. / lapackf77_dlapy2( &d__1, &d__2 );
blasf77_dscal( &n, &scl, VR(0,i), &ione );
blasf77_dscal( &n, &scl, VR(0,i+1), &ione );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = *VR(k,i);
d__2 = *VR(k,i+1);
work[iwrk + k] = d__1*d__1 + d__2*d__2;
}
k = blasf77_idamax( &n, &work[iwrk], &ione ) - 1; // subtract 1; k is 0-based
lapackf77_dlartg( VR(k,i), VR(k,i+1), &cs, &sn, &r );
blasf77_drot( &n, VR(0,i), &ione, VR(0,i+1), &ione, &cs, &sn );
*VR(k,i+1) = 0.;
}
}
}
CLEANUP:
/* Undo scaling if necessary */
if (scalea) {
// converged eigenvalues, stored in wr[i+1:n] and wi[i+1:n] for i = INFO
magma_int_t nval = n - (*info);
magma_int_t ld = max( nval, 1 );
lapackf77_dlascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wr + (*info), &ld, &ierr );
lapackf77_dlascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wi + (*info), &ld, &ierr );
if (*info > 0) {
// first ilo columns were already upper triangular,
// so the corresponding eigenvalues are also valid.
nval = ilo - 1;
lapackf77_dlascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wr, &n, &ierr );
lapackf77_dlascl( "G", &izero, &izero, &cscale, &anrm, &nval, &ione, wi, &n, &ierr );
}
}
#if defined(Version3) || defined(Version4) || defined(Version5)
magma_free( dT );
#endif
#if defined(Version4) || defined(Version5)
magma_free_cpu( T );
#endif
timer_stop( time_total );
flops_stop( flop_total );
timer_printf( "dgeev times n %5d, gehrd %7.3f, unghr %7.3f, hseqr %7.3f, trevc %7.3f, total %7.3f, sum %7.3f\n",
(int) n, time_gehrd, time_unghr, time_hseqr, time_trevc, time_total, time_sum );
timer_printf( "dgeev flops n %5d, gehrd %7lld, unghr %7lld, hseqr %7lld, trevc %7lld, total %7lld, sum %7lld\n",
(int) n, flop_gehrd, flop_unghr, flop_hseqr, flop_trevc, flop_total, flop_sum );
work[0] = MAGMA_D_MAKE( (double) optwrk, 0. );
return *info;
} /* magma_dgeev */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dlaset
Code is very similar to testing_dlacpy.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_R;
magmaDouble_ptr d_A;
double offdiag = MAGMA_D_MAKE( 1.2000, 6.7000 );
double diag = MAGMA_D_MAKE( 3.1415, 2.7183 );
magma_int_t M, N, size, lda, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("=================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
//M += 2; // space for insets
//N += 2;
lda = M;
ldda = ((M+31)/32)*32;
size = lda*N;
if ( uplo[iuplo] == MagmaLower || uplo[iuplo] == MagmaUpper ) {
// save triangle (with diagonal)
// TODO wrong for trapezoid
gbytes = sizeof(double) * 0.5*N*(N+1) / 1e9;
}
else {
// save entire matrix
gbytes = sizeof(double) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, double, size );
TESTING_MALLOC_CPU( h_R, double, size );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_D_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix( M, N, h_A, lda, d_A, 0, ldda, opts.queue );
gpu_time = magma_sync_wtime( 0 );
//magmablas_dlaset( uplo[iuplo], M-2, N-2, offdiag, diag, d_A+1+ldda, 0, ldda, opts.queue ); // inset by 1 row & col
magmablas_dlaset( uplo[iuplo], M, N, offdiag, diag, d_A, 0, ldda, opts.queue );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//magma_int_t M2 = M-2; // inset by 1 row & col
//magma_int_t N2 = N-2;
//lapackf77_dlaset( lapack_uplo_const( uplo[iuplo] ), &M2, &N2, &offdiag, &diag, h_A+1+lda, &lda );
lapackf77_dlaset( lapack_uplo_const( uplo[iuplo] ), &M, &N, &offdiag, &diag, h_A, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if ( opts.verbose ) {
printf( "A= " ); magma_dprint( M, N, h_A, lda );
printf( "dA=" ); magma_dprint_gpu( M, N, d_A, 0, ldda, opts.queue );
}
/* =====================================================================
Check the result
=================================================================== */
magma_dgetmatrix( M, N, d_A, 0, ldda, h_R, lda, opts.queue );
blasf77_daxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work);
printf("%5s %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
lapack_uplo_const( uplo[iuplo] ), (int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgesv_gpu
*/
int main(int argc , char **argv)
{
real_Double_t gflops, gpu_perf, gpu_time;
double Rnorm, Anorm, Xnorm, *work;
double *hA, *hB, *hX;
magmaDouble_ptr dA, dB;
magma_int_t *ipiv;
magma_int_t N = 0, n2, lda, ldb, ldda, lddb;
magma_int_t size[7] =
{ 1024, 2048, 3072, 4032, 5184, 6048, 7000};
magma_int_t i, info, szeB;
double z_one = MAGMA_D_ONE;
double mz_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t NRHS = 100;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
if (strcmp("-R", argv[i])==0)
NRHS = atoi(argv[++i]);
}
if (N>0) size[0] = size[6] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_dgesv_gpu -N <matrix size> -R <right hand sides>\n\n");
}
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
/* Allocate memory for the largest matrix */
N = size[6];
n2 = N * N;
ldda = ((N+31)/32) * 32;
// ldda = N;
lddb = ldda;
TESTING_MALLOC_PIN( ipiv, magma_int_t, N );
TESTING_MALLOC_PIN( hA, double, n2 );
TESTING_MALLOC_PIN( hB, double, N*NRHS );
TESTING_MALLOC_PIN( hX, double, N*NRHS );
TESTING_MALLOC_PIN( work, double, N );
TESTING_MALLOC_DEV( dA, double, ldda*N );
TESTING_MALLOC_DEV( dB, double, lddb*NRHS );
printf("\n\n");
printf(" N NRHS GPU GFlop/s (sec) ||B - AX|| / ||A||*||X||\n");
printf("===========================================================\n");
for( i = 0; i < 7; i++ ) {
N = size[i];
lda = N;
ldb = lda;
n2 = lda*N;
szeB = ldb*NRHS;
ldda = ((N+31)/32)*32;
//ldda = N;
lddb = ldda;
gflops = ( FLOPS_GETRF( (double)N, (double)N ) +
FLOPS_GETRS( (double)N, (double)NRHS ) ) / 1e9;
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &n2, hA );
lapackf77_dlarnv( &ione, ISEED, &szeB, hB );
/* Warm up to measure the performance */
magma_dsetmatrix( N, N, hA, 0, lda, dA, 0, ldda, queue );
magma_dsetmatrix( N, NRHS, hB, 0, lda, dB, 0, lddb, queue );
magma_dgesv_gpu( N, NRHS, dA, 0, ldda, ipiv, dB, 0, lddb, &info, queue );
//=====================================================================
// Solve Ax = b through an LU factorization
//=====================================================================
magma_dsetmatrix( N, N, hA, 0, lda, dA, 0, ldda, queue );
magma_dsetmatrix( N, NRHS, hB, 0, lda, dB, 0, lddb, queue );
gpu_time = magma_wtime();
magma_dgesv_gpu( N, NRHS, dA, 0, ldda, ipiv, dB, 0, lddb, &info, queue );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf( "magma_dposv had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* =====================================================================
Residual
=================================================================== */
magma_dgetmatrix( N, NRHS, dB, 0, lddb, hX, 0, ldb, queue );
Anorm = lapackf77_dlange("I", &N, &N, hA, &lda, work);
Xnorm = lapackf77_dlange("I", &N, &NRHS, hX, &ldb, work);
blasf77_dgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &NRHS, &N,
&z_one, hA, &lda,
hX, &ldb,
&mz_one, hB, &ldb );
Rnorm = lapackf77_dlange("I", &N, &NRHS, hB, &ldb, work);
printf( "%5d %5d %7.2f (%7.2f) %8.2e\n",
N, NRHS, gpu_perf, gpu_time, Rnorm/(Anorm*Xnorm) );
if (argc != 1)
break;
}
/* clean up */
TESTING_FREE_PIN( hA );
TESTING_FREE_PIN( hB );
TESTING_FREE_PIN( hX );
TESTING_FREE_PIN( work );
TESTING_FREE_PIN( ipiv );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dB );
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_R, *tau, *h_work, tmp[1];
magma_int_t M, N, n2, lda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol, eps = lapackf77_dlamch("E");
tol = opts.tolerance * eps;
opts.lapack |= ( opts.check == 2 ); // check (-c2) implies lapack (-l)
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("===============================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
}
for( int 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_dgeqrf_nb(M);
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lwork = max( lwork, max( N*nb, 2*nb*nb ));
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_work, double, lwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
/* Initialize the matrix */
for ( int j=0; j<4; j++ )
ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
if ( opts.warmup ) {
magma_dgeqrf( M, N, h_R, lda, tau, h_work, lwork, opts.queues2, &info );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dgeqrf( M, N, h_R, lda, tau, h_work, lwork, opts.queues2, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
double *tau;
TESTING_MALLOC_CPU( tau, double, min_mn );
cpu_time = magma_wtime();
lapackf77_dgeqrf(&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_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau );
}
if ( opts.check == 1 && M >= N ) {
/* =====================================================================
Check the result -- dqrt02 requires M >= N
=================================================================== */
magma_int_t lwork = n2+N;
double *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC_CPU( h_W1, double, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, double, n2 ); // R
TESTING_MALLOC_CPU( h_W3, double, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, double, M ); // RWORK
lapackf77_dlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_dqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0], results[1] );
} 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 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("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 < 2 * tol ? "ok" : "failed"));
status += ! (error < 2 * 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 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_dgeev(
char jobvl, char jobvr, magma_int_t n,
double *A, magma_int_t lda,
double *WR, double *WI,
double *vl, magma_int_t ldvl,
double *vr, magma_int_t ldvr,
double *work, magma_int_t lwork,
magma_int_t *info )
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
DGEEV computes for an N-by-N real nonsymmetric matrix A, the
eigenvalues and, optionally, the left and/or right eigenvectors.
The right eigenvector v(j) of A satisfies
A * v(j) = lambda(j) * v(j)
where lambda(j) is its eigenvalue.
The left eigenvector u(j) of A satisfies
u(j)**T * A = lambda(j) * u(j)**T
where u(j)**T denotes the transpose of u(j).
The computed eigenvectors are normalized to have Euclidean norm
equal to 1 and largest component real.
Arguments
=========
JOBVL (input) CHARACTER*1
= 'N': left eigenvectors of A are not computed;
= 'V': left eigenvectors of are computed.
JOBVR (input) CHARACTER*1
= 'N': right eigenvectors of A are not computed;
= 'V': right eigenvectors of A are computed.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) DOUBLE PRECISION array, dimension (LDA,N)
On entry, the N-by-N matrix A.
On exit, A has been overwritten.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
WR (output) DOUBLE PRECISION array, dimension (N)
WI (output) DOUBLE PRECISION array, dimension (N)
WR and WI contain the real and imaginary parts,
respectively, of the computed eigenvalues. Complex
conjugate pairs of eigenvalues appear consecutively
with the eigenvalue having the positive imaginary part
first.
VL (output) DOUBLE PRECISION array, dimension (LDVL,N)
If JOBVL = 'V', the left eigenvectors u(j) are stored one
after another in the columns of VL, in the same order
as their eigenvalues.
If JOBVL = 'N', VL is not referenced.
u(j) = VL(:,j), the j-th column of VL.
LDVL (input) INTEGER
The leading dimension of the array VL. LDVL >= 1; if
JOBVL = 'V', LDVL >= N.
VR (output) DOUBLE PRECISION array, dimension (LDVR,N)
If JOBVR = 'V', the right eigenvectors v(j) are stored one
after another in the columns of VR, in the same order
as their eigenvalues.
If JOBVR = 'N', VR is not referenced.
v(j) = VR(:,j), the j-th column of VR.
LDVR (input) INTEGER
The leading dimension of the array VR. LDVR >= 1; if
JOBVR = 'V', LDVR >= N.
WORK (workspace/output) DOUBLE PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= (2+nb)*N.
For optimal performance, LWORK >= (2+2*nb)*N.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value.
> 0: if INFO = i, the QR algorithm failed to compute all the
eigenvalues, and no eigenvectors have been computed;
elements and i+1:N of W contain eigenvalues which have
converged.
===================================================================== */
#define vl(i,j) (vl + (i) + (j)*ldvl)
#define vr(i,j) (vr + (i) + (j)*ldvr)
magma_int_t c_one = 1;
magma_int_t c_zero = 0;
double d__1, d__2;
double r, cs, sn, scl;
double dum[1], eps;
double anrm, cscale, bignum, smlnum;
magma_int_t i, k, ilo, ihi;
magma_int_t ibal, ierr, itau, iwrk, nout, liwrk, i__1, i__2, nb;
magma_int_t scalea, minwrk, lquery, wantvl, wantvr, select[1];
char side[2] = {0, 0};
char jobvl_[2] = {jobvl, 0};
char jobvr_[2] = {jobvr, 0};
*info = 0;
lquery = lwork == -1;
wantvl = lapackf77_lsame( jobvl_, "V" );
wantvr = lapackf77_lsame( jobvr_, "V" );
if (! wantvl && ! lapackf77_lsame( jobvl_, "N" )) {
*info = -1;
} else if (! wantvr && ! lapackf77_lsame( jobvr_, "N" )) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if ( (ldvl < 1) || (wantvl && (ldvl < n))) {
*info = -9;
} else if ( (ldvr < 1) || (wantvr && (ldvr < n))) {
*info = -11;
}
/* Compute workspace */
nb = magma_get_dgehrd_nb( n );
if (*info == 0) {
minwrk = (2+nb)*n;
work[0] = MAGMA_D_MAKE( (double) minwrk, 0. );
if (lwork < minwrk && ! lquery) {
*info = -13;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (n == 0) {
return *info;
}
#if defined(VERSION3)
double *dT;
if (MAGMA_SUCCESS != magma_dmalloc( &dT, nb*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
#endif
/* Get machine constants */
eps = lapackf77_dlamch( "P" );
smlnum = lapackf77_dlamch( "S" );
bignum = 1. / smlnum;
lapackf77_dlabad( &smlnum, &bignum );
smlnum = magma_dsqrt( smlnum ) / eps;
bignum = 1. / smlnum;
/* Scale A if max element outside range [SMLNUM,BIGNUM] */
anrm = lapackf77_dlange( "M", &n, &n, A, &lda, dum );
scalea = 0;
if (anrm > 0. && anrm < smlnum) {
scalea = 1;
cscale = smlnum;
} else if (anrm > bignum) {
scalea = 1;
cscale = bignum;
}
if (scalea) {
lapackf77_dlascl( "G", &c_zero, &c_zero, &anrm, &cscale, &n, &n, A, &lda, &ierr );
}
/* Balance the matrix
* (Workspace: need N) */
ibal = 0;
lapackf77_dgebal( "B", &n, A, &lda, &ilo, &ihi, &work[ibal], &ierr );
/* Reduce to upper Hessenberg form
* (Workspace: need 3*N, prefer 2*N + N*NB) */
itau = ibal + n;
iwrk = itau + n;
liwrk = lwork - iwrk;
#if defined(VERSION1)
// Version 1 - LAPACK
lapackf77_dgehrd( &n, &ilo, &ihi, A, &lda,
&work[itau], &work[iwrk], &liwrk, &ierr );
#elif defined(VERSION2)
// Version 2 - LAPACK consistent HRD
magma_dgehrd2( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, &ierr );
#elif defined(VERSION3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored,
magma_dgehrd( n, ilo, ihi, A, lda,
&work[itau], &work[iwrk], liwrk, dT, &ierr );
#endif
if (wantvl) {
/* Want left eigenvectors
* Copy Householder vectors to VL */
side[0] = 'L';
lapackf77_dlacpy( MagmaLowerStr, &n, &n,
A, &lda, vl, &ldvl );
/* Generate orthogonal matrix in VL
* (Workspace: need 3*N-1, prefer 2*N + (N-1)*NB) */
#if defined(VERSION1) || defined(VERSION2)
// Version 1 & 2 - LAPACK
lapackf77_dorghr( &n, &ilo, &ihi, vl, &ldvl, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(VERSION3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_dorghr( n, ilo, ihi, vl, ldvl, &work[itau], dT, nb, &ierr );
#endif
/* Perform QR iteration, accumulating Schur vectors in VL
* (Workspace: need N+1, prefer N+HSWORK (see comments) ) */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_dhseqr( "S", "V", &n, &ilo, &ihi, A, &lda, WR, WI,
vl, &ldvl, &work[iwrk], &liwrk, info );
if (wantvr) {
/* Want left and right eigenvectors
* Copy Schur vectors to VR */
side[0] = 'B';
lapackf77_dlacpy( "F", &n, &n, vl, &ldvl, vr, &ldvr );
}
}
else if (wantvr) {
/* Want right eigenvectors
* Copy Householder vectors to VR */
side[0] = 'R';
lapackf77_dlacpy( "L", &n, &n, A, &lda, vr, &ldvr );
/* Generate orthogonal matrix in VR
* (Workspace: need 3*N-1, prefer 2*N + (N-1)*NB) */
#if defined(VERSION1) || defined(VERSION2)
// Version 1 & 2 - LAPACK
lapackf77_dorghr( &n, &ilo, &ihi, vr, &ldvr, &work[itau],
&work[iwrk], &liwrk, &ierr );
#elif defined(VERSION3)
// Version 3 - LAPACK consistent MAGMA HRD + T matrices stored
magma_dorghr( n, ilo, ihi, vr, ldvr, &work[itau], dT, nb, &ierr );
#endif
/* Perform QR iteration, accumulating Schur vectors in VR
* (Workspace: need N+1, prefer N+HSWORK (see comments) ) */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_dhseqr( "S", "V", &n, &ilo, &ihi, A, &lda, WR, WI,
vr, &ldvr, &work[iwrk], &liwrk, info );
}
else {
/* Compute eigenvalues only
* (Workspace: need N+1, prefer N+HSWORK (see comments) ) */
iwrk = itau;
liwrk = lwork - iwrk;
lapackf77_dhseqr( "E", "N", &n, &ilo, &ihi, A, &lda, WR, WI,
vr, &ldvr, &work[iwrk], &liwrk, info );
}
/* If INFO > 0 from DHSEQR, then quit */
if (*info > 0) {
goto CLEANUP;
}
if (wantvl || wantvr) {
/* Compute left and/or right eigenvectors
* (Workspace: need 4*N) */
liwrk = lwork - iwrk;
#if TREVC_VERSION == 1
lapackf77_dtrevc( side, "B", select, &n, A, &lda, vl, &ldvl,
vr, &ldvr, &n, &nout, &work[iwrk], &ierr );
#elif TREVC_VERSION == 2
lapackf77_dtrevc3( side, "B", select, &n, A, &lda, vl, &ldvl,
vr, &ldvr, &n, &nout, &work[iwrk], &liwrk, &ierr );
#endif
}
if (wantvl) {
/* Undo balancing of left eigenvectors
* (Workspace: need N) */
lapackf77_dgebak( "B", "L", &n, &ilo, &ihi, &work[ibal], &n,
vl, &ldvl, &ierr );
/* Normalize left eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
if ( WI[i] == 0. ) {
scl = 1. / cblas_dnrm2( n, vl(0,i), 1 );
cblas_dscal( n, scl, vl(0,i), 1 );
}
else if ( WI[i] > 0. ) {
d__1 = cblas_dnrm2( n, vl(0,i), 1 );
d__2 = cblas_dnrm2( n, vl(0,i+1), 1 );
scl = 1. / lapackf77_dlapy2( &d__1, &d__2 );
cblas_dscal( n, scl, vl(0,i), 1 );
cblas_dscal( n, scl, vl(0,i+1), 1 );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = *vl(k,i);
d__2 = *vl(k,i+1);
work[iwrk + k] = d__1*d__1 + d__2*d__2;
}
k = cblas_idamax( n, &work[iwrk], 1 );
lapackf77_dlartg( vl(k,i), vl(k,i+1), &cs, &sn, &r );
cblas_drot( n, vl(0,i), 1, vl(0,i+1), 1, cs, sn );
*vl(k,i+1) = 0.;
}
}
}
if (wantvr) {
/* Undo balancing of right eigenvectors
* (Workspace: need N) */
lapackf77_dgebak( "B", "R", &n, &ilo, &ihi, &work[ibal], &n,
vr, &ldvr, &ierr );
/* Normalize right eigenvectors and make largest component real */
for (i = 0; i < n; ++i) {
if ( WI[i] == 0. ) {
scl = 1. / cblas_dnrm2( n, vr(0,i), 1 );
cblas_dscal( n, scl, vr(0,i), 1 );
}
else if ( WI[i] > 0. ) {
d__1 = cblas_dnrm2( n, vr(0,i), 1 );
d__2 = cblas_dnrm2( n, vr(0,i+1), 1 );
scl = 1. / lapackf77_dlapy2( &d__1, &d__2 );
cblas_dscal( n, scl, vr(0,i), 1 );
cblas_dscal( n, scl, vr(0,i+1), 1 );
for (k = 0; k < n; ++k) {
/* Computing 2nd power */
d__1 = *vr(k,i);
d__2 = *vr(k,i+1);
work[iwrk + k] = d__1*d__1 + d__2*d__2;
}
k = cblas_idamax( n, &work[iwrk], 1 );
lapackf77_dlartg( vr(k,i), vr(k,i+1), &cs, &sn, &r );
cblas_drot( n, vr(0,i), 1, vr(0,i+1), 1, cs, sn );
*vr(k,i+1) = 0.;
}
}
}
CLEANUP:
/* Undo scaling if necessary */
if (scalea) {
i__1 = n - (*info);
i__2 = max( n - (*info), 1 );
lapackf77_dlascl( "G", &c_zero, &c_zero, &cscale, &anrm, &i__1, &c_one,
WR + (*info), &i__2, &ierr );
lapackf77_dlascl( "G", &c_zero, &c_zero, &cscale, &anrm, &i__1, &c_one,
WI + (*info), &i__2, &ierr );
if (*info > 0) {
i__1 = ilo - 1;
lapackf77_dlascl( "G", &c_zero, &c_zero, &cscale, &anrm, &i__1, &c_one,
WR, &n, &ierr );
lapackf77_dlascl( "G", &c_zero, &c_zero, &cscale, &anrm, &i__1, &c_one,
WI, &n, &ierr );
}
}
#if defined(VERSION3)
magma_free( dT );
#endif
return *info;
} /* magma_dgeev */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqlf
*/
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_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_dgeqrf_nb(M);
gflops = FLOPS_DGEQLF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgeqlf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lwork = max( lwork, N*nb );
lwork = max( lwork, 2*nb*nb);
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_work, double, lwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dgeqlf( 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_dgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_dgeqlf(&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_dgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("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;
}
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double *h_x, *h_x2, *h_tau, *h_tau2;
magmaDouble_ptr d_x, d_tau;
double c_neg_one = MAGMA_D_NEG_ONE;
double error, error2, work[1];
magma_int_t N, nb, lda, ldda, size;
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");
// does larfg on nb columns, one after another
nb = (opts.nb > 0 ? opts.nb : 64);
magma_queue_t queue = 0;
printf(" N nb CPU GFLop/s (ms) GPU GFlop/s (ms) error tau error\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 = ((N+31)/32)*32;
gflops = FLOPS_DLARFG( N ) / 1e9 * nb;
TESTING_MALLOC_CPU( h_x, double, N*nb );
TESTING_MALLOC_CPU( h_x2, double, N*nb );
TESTING_MALLOC_CPU( h_tau, double, nb );
TESTING_MALLOC_CPU( h_tau2, double, nb );
TESTING_MALLOC_DEV( d_x, double, ldda*nb );
TESTING_MALLOC_DEV( d_tau, double, nb );
/* Initialize the vectors */
size = N*nb;
lapackf77_dlarnv( &ione, ISEED, &size, h_x );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_dsetmatrix( N, nb, h_x, N, d_x, ldda );
gpu_time = magma_sync_wtime( queue );
for( int j = 0; j < nb; ++j ) {
magmablas_dlarfg( N, &d_x[0+j*ldda], &d_x[1+j*ldda], ione, &d_tau[j] );
}
gpu_time = magma_sync_wtime( queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_dgetmatrix( N, nb, d_x, ldda, h_x2, N );
magma_dgetvector( nb, d_tau, 1, h_tau2, 1 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for( int j = 0; j < nb; ++j ) {
lapackf77_dlarfg( &N, &h_x[0+j*lda], &h_x[1+j*lda], &ione, &h_tau[j] );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Error Computation and Performance Comparison
=================================================================== */
blasf77_daxpy( &size, &c_neg_one, h_x, &ione, h_x2, &ione );
error = lapackf77_dlange( "F", &N, &nb, h_x2, &N, work )
/ lapackf77_dlange( "F", &N, &nb, h_x, &N, work );
// tau can be 0
blasf77_daxpy( &nb, &c_neg_one, h_tau, &ione, h_tau2, &ione );
error2 = lapackf77_dlange( "F", &nb, &ione, h_tau, &nb, work );
if ( error2 != 0 ) {
error2 = lapackf77_dlange( "F", &nb, &ione, h_tau2, &nb, work ) / error2;
}
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) N, (int) nb, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, error2,
(error < tol && error2 < tol ? "ok" : "failed") );
status += ! (error < tol && error2 < tol);
TESTING_FREE_CPU( h_x );
TESTING_FREE_CPU( h_x2 );
TESTING_FREE_CPU( h_tau );
TESTING_FREE_DEV( d_x );
TESTING_FREE_DEV( d_tau );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
DSGESV computes the solution to a real system of linear equations
A * X = B, A**T * X = B, or A**H * X = B,
where A is an N-by-N matrix and X and B are N-by-NRHS matrices.
DSGESV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (ldda,N)
On entry, the N-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
@param[out]
ipiv INTEGER array, dimension (N)
The pivot indices that define the permutation matrix P;
row i of the matrix was interchanged with row IPIV(i).
Corresponds either to the single precision factorization
(if info.EQ.0 and ITER.GE.0) or the double precision
factorization (if info.EQ.0 and ITER.LT.0).
@param[out]
dipiv INTEGER array on the GPU, dimension (N)
The pivot indices; for 1 <= i <= N, after permuting, row i of the
matrix was moved to row dIPIV(i).
Note this is different than IPIV, where interchanges
are applied one-after-another.
@param[in]
dB DOUBLE PRECISION array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
@param[in]
lddb INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
@param[out]
dX DOUBLE PRECISION array on the GPU, dimension (lddx,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
@param
dworkd (workspace) DOUBLE PRECISION array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dworks (workspace) SINGLE PRECISION array on the GPU, dimension (N*(N+NRHS))
This array is used to store the real single precision matrix
and the right-hand sides or solutions in single precision.
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, double precision
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SGETRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
- > 0: if info = i, U(i,i) computed in DOUBLE PRECISION is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
@ingroup magma_dgesv_driver
********************************************************************/
extern "C" magma_int_t
magma_dsgesv_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
double *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *dipiv,
double *dB, magma_int_t lddb,
double *dX, magma_int_t lddx,
double *dworkd, float *dworks,
magma_int_t *iter, magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
double *dR;
float *dSA, *dSX;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlange(MagmaInfNorm, n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
//magmablas_dlag2s( n, nrhs, dB, lddb, dSX, lddsx, info ); // done inside dsgetrs with pivots
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_dlag2s( n, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_sgetrf_gpu( n, n, dSA, lddsa, ipiv, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Generate parallel pivots
{
magma_int_t *newipiv;
magma_imalloc_cpu( &newipiv, n );
if ( newipiv == NULL ) {
*iter = -3;
goto FALLBACK;
}
swp2pswp( trans, n, ipiv, newipiv );
magma_setvector( n, sizeof(magma_int_t), newipiv, 1, dipiv, 1 );
magma_free_cpu( newipiv );
}
// solve dSA*dSX = dB in single precision
// converts dB to dSX and applies pivots, solves, then converts result back to dX
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dB, lddb, dX, lddx, dSX, info );
// residual dR = dB - dA*dX in double precision
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
// solve dSA*dSX = R in single precision
// convert result back to double precision dR
// it's okay that dR is used for both dB input and dX output.
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dR, lddr, dR, lddr, dSX, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_daxpycp( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_dgetrf_gpu( n, n, dA, ldda, ipiv, info );
if (*info == 0) {
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_dgetrs_gpu( trans, n, nrhs, dA, ldda, ipiv, dX, lddx, info );
}
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_dsymm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 3.456, 5.678 );
double beta = MAGMA_D_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double Anorm, error, work[1];
double *hA, *hB, *hC, *hR;
magmaDouble_ptr dA[MagmaMaxGPUs], dB[MagmaMaxGPUs], dC[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magmaDouble_ptr dA2;
magma_int_t i, j, dev, M, N, size, lda, ldb, ldc, ldda, lddb, lddc, msize, nb;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t ncmplx = 0;
magma_buildconnection_mgpu( gnode, &ncmplx, opts.ngpu );
printf("%% Initializing communication pattern... GPU-ncmplx %d\n", (int) ncmplx);
for (i=0; i < ncmplx; ++i) {
magma_int_t myngpu = gnode[i][MagmaMaxGPUs];
printf("%% cmplx %d has %d GPUs:", i, myngpu);
for (j=0; j < myngpu; ++j) {
printf(" %d", (int) gnode[i][j]);
if (j < myngpu-1) {
printf(",");
}
}
printf("\n");
}
// number of queues per GPU. Requires ngpu.
magma_int_t nqueue = opts.ngpu;
// number of events per GPU. Require ngpu*ngpu.
magma_int_t nevents = opts.ngpu*opts.ngpu;
magma_queue_t queues[MagmaMaxGPUs][20], queues0[MagmaMaxGPUs];
magma_event_t events[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs + 10];
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( i = 0; i < nqueue; ++i ) {
magma_queue_create( dev, &queues[dev][i] );
}
queues0[dev] = queues[dev][0];
for( i = 0; i < nevents; ++i ) {
cudaEventCreateWithFlags( &events[dev][i], cudaEventDisableTiming );
}
}
printf("%% nb %d, ngpu %d, version %d\n", (int) nb, (int) opts.ngpu, (int) opts.version );
printf("%% M N nb offset CPU Gflop/s (sec) GPU Gflop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||B||\n");
printf("%%========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M; // TODO depends on side
ldb = M;
ldc = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
lddc = magma_roundup( ldc, opts.align ); // multiple of 32 by default
gflops = FLOPS_DSYMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = lddc*N + (M*N)*opts.ngpu;
TESTING_MALLOC_CPU( hA, double, lda*M );
TESTING_MALLOC_CPU( hB, double, ldb*N );
TESTING_MALLOC_CPU( hC, double, ldc*N );
TESTING_MALLOC_PIN( hR, double, ldc*N );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( dA[dev], double, ldda*mlocal );
TESTING_MALLOC_DEV( dB[dev], double, lddb*N );
TESTING_MALLOC_DEV( dC[dev], double, lddc*N );
TESTING_MALLOC_DEV( dwork[dev], double, dworksiz );
}
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, double, ldda*M );
}
size = lda*M;
lapackf77_dlarnv( &ione, iseed, &size, hA );
magma_dmake_symmetric( M, hA, lda );
size = ldb*N;
lapackf77_dlarnv( &ione, iseed, &size, hB );
size = ldc*N;
lapackf77_dlarnv( &ione, iseed, &size, hC );
lapackf77_dlacpy( "Full", &M, &N, hC, &ldc, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb, queues0 );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_dsetmatrix( M, N, hB, lda, dB[dev], ldda, opts.queue );
// since when offset != 0, the GPU that does beta*C may not be 0,
// send initial hC to all GPUs.
magma_dsetmatrix( M, N, hC, lda, dC[dev], ldda, opts.queue );
}
trace_init( 1, opts.ngpu, nqueue, (magma_queue_t*) queues );
gpu_time = magma_sync_wtime(0);
magmablas_dsymm_mgpu(
MagmaLeft, MagmaLower, msize, N,
alpha, dA, ldda, offset,
dB, ldda,
beta, dC, ldda, dwork, dworksiz,
opts.ngpu, nb, queues, nqueue, events, nevents, gnode, ncmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "dsymm-m%d-n%d-nb%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magma_dsetmatrix( M, M, hA, lda, dA2, ldda, opts.queue );
magma_dsetmatrix( M, N, hB, lda, dB[0], ldda, opts.queue );
magma_dsetmatrix( M, N, hC, lda, dwork[0], ldda, opts.queue );
gpu_time2 = magma_sync_wtime(0);
magma_dsymm(
MagmaLeft, MagmaLower, msize, N,
alpha, dA2 + offset + offset*ldda, ldda,
dB[0], ldda,
beta, dwork[0], ldda, opts.queue );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||B||
Anorm = lapackf77_dlange("fro", &msize, &msize, hA + offset + offset*lda, &lda, work );
Anorm *= lapackf77_dlange("fro", &msize, &N, hB, &lda, work );
//printf( "A =" ); magma_dprint( M, M, hA, lda );
//printf( "B =" ); magma_dprint( M, N, hB, lda );
//printf( "C =" ); magma_dprint( M, N, hC, lda );
cpu_time = magma_wtime();
blasf77_dsymm( "Left", "Lower", &msize, &N,
&alpha, hA + offset + offset*lda, &lda,
hB, &lda,
&beta, hC, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
for (dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_dgetmatrix( M, N, dC[dev], ldda, hR, lda, opts.queue );
// compute relative error ||R||/||A||*||B||, where R := C_magma - C_lapack = R - C
size = ldc*N;
blasf77_daxpy( &size, &c_neg_one, hC, &ione, hR, &ione );
error = lapackf77_dlange("fro", &msize, &N, hR, &lda, work) / Anorm;
//printf( "R =" ); magma_dprint( M, N, hR, lda );
bool okay = (error < tol);
status += ! okay;
if (dev == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (okay ? "ok" : "failed") );
}
else {
printf( " dev %d %74s %8.2e %s\n", dev, "",
error, (okay ? "ok" : "failed") );
}
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hB );
TESTING_FREE_CPU( hC );
TESTING_FREE_PIN( hR );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
TESTING_FREE_DEV( dA[dev] );
TESTING_FREE_DEV( dB[dev] );
TESTING_FREE_DEV( dC[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( i = 0; i < nqueue; ++i ) {
magma_queue_destroy( queues[dev][i] );
}
for( i = 0; i < nevents; ++i ) {
magma_event_destroy( events[dev][i] );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
