int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
float error, work[1];
int transA = MagmaNoTrans;
int transB = MagmaNoTrans;
float Cnorm;
magma_int_t istart = 1024;
magma_int_t iend = 8194;
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};
magmaFloatComplex *h_A, *h_B, *h_C, *h_C2;
magmaFloatComplex_ptr d_A, d_B, d_C;
magmaFloatComplex mzone = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.48, 0.38 );
if (argc != 1){
for(i=1; i<argc; i++){
if ( strcmp("-N", argv[i]) == 0 ){
N0 = atoi(argv[++i]);
}
else if ( strcmp("-M", argv[i]) == 0 ){
M0 = atoi(argv[++i]);
}
else if ( strcmp("-K", argv[i]) == 0 ){
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;
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else if (strcmp("-NC", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaConjTrans;
}
else if (strcmp("-TC", argv[i])==0){
transA = MagmaTrans;
transB = MagmaConjTrans;
}
else if (strcmp("-CN", argv[i])==0){
transA = MagmaConjTrans;
transB = MagmaNoTrans;
}
else if (strcmp("-CT", argv[i])==0){
transA = MagmaConjTrans;
transB = MagmaTrans;
}
else if (strcmp("-CC", argv[i])==0){
transA = transB = MagmaConjTrans;
}
#endif
}
}
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 = lddc = ((M+31)/32)*32;
lddb = ((ldb+31)/32)*32;
K+=32;
M+=32;
N +=32;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*K );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*Bn );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_C2, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*K );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N );
printf("\nUsage: \n");
printf(" testing_cgemm [-NN|NT|TN|TT] [-N %d] \n\n", 1024);
printf("\n");
printf("Testing transA = %c transB = %c\n", transA, transB);
printf(" M N K clAmdBlas GFLop/s (sec) CPU GFlop/s (sec) error\n");
printf("===========================================================================\n");
for(i=istart; i<iend; i = (int)(i*1.25) )
{
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( (float)M, (float)N, (float)K ) * 1e-9;
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_clarnv( &ione, ISEED, &szeA, h_A );
lapackf77_clarnv( &ione, ISEED, &szeB, h_B );
lapackf77_clarnv( &ione, ISEED, &szeC, h_C );
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_csetmatrix( Am, An, h_A, 0, lda, d_A, 0, ldda, queue );
magma_csetmatrix( Bm, Bn, h_B, 0, ldb, d_B, 0, lddb, queue );
magma_csetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_cgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_csetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_queue_sync( queue );
gpu_time = magma_wtime();
magma_cgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_queue_sync( queue);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
magma_cgetmatrix( M, N, d_C, 0, lddc, h_C2, 0, ldc, queue );
/* =====================================================================
Performs operation using CPU-BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_cgemm( 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;
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_clange( "M", &M, &N, h_C, &ldc, work );
/* =====================================================================
Error Computation and Performance Compariosn
=================================================================== */
blasf77_caxpy(&szeC, &mzone, h_C, &ione, h_C2, &ione);
error = lapackf77_clange("M", &M, &N, h_C2, &ldc, work)/Cnorm;
printf("%5d %5d %5d %8.2f (%6.2f) %6.2f (%6.2f) %e\n",
M, N, K, gpu_perf, gpu_time, cpu_perf, cpu_time, error);
}
/* 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_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zherk
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
double cublas_error, Cnorm, work[1];
magma_int_t N, K;
magma_int_t Ak, An;
magma_int_t sizeA, sizeC;
magma_int_t lda, ldc, ldda, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex *h_A, *h_C, *h_Ccublas;
magmaDoubleComplex *d_A, *d_C;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
double beta = MAGMA_D_MAKE( -0.48, 0.38 );
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("If running lapack (option --lapack), CUBLAS error is computed\n"
"relative to CPU BLAS result.\n\n");
printf("uplo = %s, transA = %s\n",
lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA) );
printf(" N K 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 ) {
N = opts.nsize[itest];
K = opts.ksize[itest];
gflops = FLOPS_ZHERK(K, N) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = An = N;
Ak = K;
} else {
lda = An = K;
Ak = N;
}
ldc = N;
ldda = ((lda+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*Ak;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC_CPU( h_C, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_CPU( h_Ccublas, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_C, magmaDoubleComplex, lddc*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( An, Ak, h_A, lda, d_A, ldda );
magma_zsetmatrix( N, N, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasZherk( handle, cublas_uplo_const(opts.uplo), cublas_trans_const(opts.transA), N, K,
&alpha, d_A, ldda,
&beta, d_C, lddc );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( N, N, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_zherk( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), &N, &K,
&alpha, h_A, &lda,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_zlanhe("fro", lapack_uplo_const(opts.uplo), &N, h_C, &ldc, work);
blasf77_zaxpy( &sizeC, &c_neg_one, h_C, &ione, h_Ccublas, &ione );
cublas_error = lapackf77_zlanhe( "fro", lapack_uplo_const(opts.uplo), &N, h_Ccublas, &ldc, work ) / Cnorm;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) K,
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) N, (int) K,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Ccublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgehrd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A, *h_R, *h_Q, *h_work, *tau, *twork, *dT;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
#endif
float eps, result[2];
magma_int_t N, n2, lda, nb, lwork, ltwork, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
eps = lapackf77_slamch( "E" );
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) |A-QHQ'|/N|A| |I-QQ'|/N\n");
printf("=========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
nb = magma_get_cgehrd_nb(N);
/* We suppose the magma nb is bigger than lapack nb */
lwork = N*nb;
gflops = FLOPS_CGEHRD( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, N );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, nb*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgehrd( N, ione, N, h_R, lda, tau, h_work, lwork, dT, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgehrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.check ) {
ltwork = 2*(N*N);
TESTING_MALLOC_PIN( h_Q, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( twork, magmaFloatComplex, ltwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_CPU( rwork, float, N );
#endif
lapackf77_clacpy(MagmaUpperLowerStr, &N, &N, h_R, &lda, h_Q, &lda);
for( int j = 0; j < N-1; ++j )
for( int i = j+2; i < N; ++i )
h_R[i+j*lda] = MAGMA_C_ZERO;
magma_cunghr(N, ione, N, h_Q, lda, tau, dT, nb, &info);
if (info != 0) {
printf("magma_cunghr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
exit(1);
}
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_chst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, rwork, result);
#else
lapackf77_chst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, result);
#endif
TESTING_FREE_PIN( h_Q );
TESTING_FREE_CPU( twork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_CPU( rwork );
#endif
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgehrd(&N, &ione, &N, h_R, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgehrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Print performance and error.
=================================================================== */
if ( opts.lapack ) {
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f)",
(int) N, gpu_perf, gpu_time );
}
if ( opts.check ) {
printf(" %8.2e %8.2e %s\n",
result[0]*eps, result[1]*eps,
( ( (result[0]*eps < tol) && (result[1]*eps < tol) ) ? "ok" : "failed") );
status += ! (result[0]*eps < tol);
status += ! (result[1]*eps < tol);
}
else {
printf(" --- ---\n");
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( tau );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double cublas_error, normA, normx, normr, work[1];
magma_int_t N, info;
magma_int_t sizeA;
magma_int_t lda, ldda;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *ipiv;
magmaDoubleComplex *h_A, *h_b, *h_x, *h_xcublas;
magmaDoubleComplex_ptr d_A, d_x;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("uplo = %s, transA = %s, diag = %s\n",
lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) );
printf(" 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 ) {
N = opts.nsize[itest];
gflops = FLOPS_ZTRSM(opts.side, N, 1) / 1e9;
lda = N;
ldda = ((lda+31)/32)*32;
sizeA = lda*N;
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( h_b, magmaDoubleComplex, N );
TESTING_MALLOC_CPU( h_x, magmaDoubleComplex, N );
TESTING_MALLOC_CPU( h_xcublas, magmaDoubleComplex, N );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( d_x, magmaDoubleComplex, N );
/* Initialize the matrices */
/* Factor A into LU to get well-conditioned triangular matrix.
* Copy L to U, since L seems okay when used with non-unit diagonal
* (i.e., from U), while U fails when used with unit diagonal. */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zgetrf( &N, &N, h_A, &lda, ipiv, &info );
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < j; ++i ) {
*h_A(i,j) = *h_A(j,i);
}
}
lapackf77_zlarnv( &ione, ISEED, &N, h_b );
blasf77_zcopy( &N, h_b, &ione, h_x, &ione );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_zsetvector( N, h_x, 1, d_x, 1 );
cublas_time = magma_sync_wtime( NULL );
cublasZtrsv( opts.handle, cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
N,
d_A, ldda,
d_x, 1 );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetvector( N, d_x, 1, h_xcublas, 1 );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ztrsv( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&N,
h_A, &lda,
h_x, &ione );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
// error for CUBLAS
normA = lapackf77_zlange( "F", &N, &N, h_A, &lda, work );
normx = lapackf77_zlange( "F", &N, &ione, h_xcublas, &ione, work );
blasf77_ztrmv( lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&N,
h_A, &lda,
h_xcublas, &ione );
blasf77_zaxpy( &N, &c_neg_one, h_b, &ione, h_xcublas, &ione );
normr = lapackf77_zlange( "F", &N, &ione, h_xcublas, &N, work );
cublas_error = normr / (normA*normx);
if ( opts.lapack ) {
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(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 %7.2f (%7.2f) --- ( --- ) %8.2e %s\n",
(int) N,
cublas_perf, 1000.*cublas_time,
cublas_error, (cublas_error < tol ? "ok" : "failed"));
status += ! (cublas_error < tol);
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_b );
TESTING_FREE_CPU( h_x );
TESTING_FREE_CPU( h_xcublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_x );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zcgeqrsv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time, gpu_perfd, gpu_perfs;
double error, gpu_error, cpu_error, Anorm, work[1];
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_A2, *h_B, *h_X, *h_R;
magmaDoubleComplex *d_A, *d_B, *d_X, *d_T;
magmaFloatComplex *d_SA, *d_SB, *d_ST;
magmaDoubleComplex *h_workd, *tau, tmp[1];
magmaFloatComplex *h_works, *tau_s;
magma_int_t lda, ldb, lhwork, lworkgpu;
magma_int_t ldda, lddb, lddx;
magma_int_t M, N, nrhs, qrsv_iters, info, size, min_mn, max_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
printf("Epsilon(double): %8.6e\n"
"Epsilon(single): %8.6e\n\n",
lapackf77_dlamch("Epsilon"), lapackf77_slamch("Epsilon") );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
nrhs = opts.nrhs;
printf(" CPU Gflop/s GPU Gflop/s |b-Ax|| / (N||A||) ||dx-x||/(N||A||)\n");
printf(" M N NRHS double double single mixed Iter CPU GPU \n");
printf("=============================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if ( M < N ) {
printf( "%5d %5d %5d skipping because M < N is not yet supported.\n", (int) M, (int) N, (int) nrhs );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
ldda = ((M+31)/32) * 32;
lddb = ((max_mn+31)/32)*32;
lddx = ((N+31)/32) * 32;
nb = max( magma_get_zgeqrf_nb( M ), magma_get_cgeqrf_nb( M ) );
gflops = (FLOPS_ZGEQRF( M, N ) + FLOPS_ZGEQRS( M, N, nrhs )) / 1e9;
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
// query for workspace size
lhwork = -1;
lapackf77_zgels( MagmaNoTransStr, &M, &N, &nrhs,
NULL, &lda, NULL, &ldb, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
lhwork = max( lhwork, lworkgpu );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( h_A2, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_workd, magmaDoubleComplex, lhwork );
tau_s = (magmaFloatComplex*)tau;
h_works = (magmaFloatComplex*)h_workd;
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaDoubleComplex, lddb*nrhs );
TESTING_MALLOC_DEV( d_X, magmaDoubleComplex, lddx*nrhs );
TESTING_MALLOC_DEV( d_T, magmaDoubleComplex, ( 2*min_mn + (N+31)/32*32 )*nb );
d_ST = (magmaFloatComplex*)d_T;
/* Initialize the matrices */
size = lda*N;
lapackf77_zlarnv( &ione, ISEED, &size, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
size = ldb*nrhs;
lapackf77_zlarnv( &ione, ISEED, &size, h_B );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
//=====================================================================
// Mixed Precision Iterative Refinement - GPU
//=====================================================================
gpu_time = magma_wtime();
magma_zcgeqrsv_gpu( M, N, nrhs,
d_A, ldda, d_B, lddb,
d_X, lddx, &qrsv_iters, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zcgeqrsv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute the residual
magma_zgetmatrix( N, nrhs, d_X, lddx, h_X, ldb );
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
Anorm = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
//=====================================================================
// Double Precision Solve
//=====================================================================
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_zgels_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_workd, lworkgpu, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perfd = gflops / gpu_time;
//=====================================================================
// Single Precision Solve
//=====================================================================
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
/* The allocation of d_SA and d_SB is done here to avoid
* to double the memory used on GPU with zcgeqrsv */
TESTING_MALLOC_DEV( d_SA, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_SB, magmaFloatComplex, lddb*nrhs );
magmablas_zlag2c( M, N, d_A, ldda, d_SA, ldda, &info );
magmablas_zlag2c( N, nrhs, d_B, lddb, d_SB, lddb, &info );
gpu_time = magma_wtime();
magma_cgels_gpu( MagmaNoTrans, M, N, nrhs, d_SA, ldda,
d_SB, lddb, h_works, lhwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perfs = gflops / gpu_time;
TESTING_FREE_DEV( d_SA );
TESTING_FREE_DEV( d_SB );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_zgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_workd, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_zgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb );
cpu_error = lapackf77_zlange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*Anorm);
gpu_error = lapackf77_zlange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
// error relative to LAPACK
size = M*nrhs;
blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_zlange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
printf("%5d %5d %5d %7.2f %7.2f %7.2f %7.2f %4d %8.2e %8.2e %8.2e %s\n",
(int) M, (int) N, (int) nrhs,
cpu_perf, gpu_perfd, gpu_perfs, gpu_perf,
(int) qrsv_iters,
cpu_error, gpu_error, error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_A2 );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( h_workd );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_X );
TESTING_FREE_DEV( d_T );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf=0., cublas_time=0., cpu_perf=0, cpu_time=0;
float error;
magma_int_t cublas_enable = 0;
magmaFloatComplex *h_A, *h_R;
magmaFloatComplex *dA_magma;
magmaFloatComplex **dA_array = NULL;
magma_int_t **dipiv_array = NULL;
magma_int_t *ipiv, *cpu_info;
magma_int_t *dipiv_magma, *dinfo_magma;
magma_int_t M, N, n2, lda, ldda, min_mn, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t batchCount;
magma_int_t status = 0;
magma_opts opts( MagmaOptsBatched );
opts.parse_opts( argc, argv );
//opts.lapack |= opts.check;
batchCount = opts.batchcount;
magma_int_t columns;
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% BatchCount M N CPU Gflop/s (ms) MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) ||PA-LU||/(||A||*N)\n");
printf("%%==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N * batchCount;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = FLOPS_CGETRF( M, N ) / 1e9 * batchCount;
TESTING_MALLOC_CPU( cpu_info, magma_int_t, batchCount );
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn * batchCount );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( dA_magma, magmaFloatComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( dipiv_magma, magma_int_t, min_mn * batchCount );
TESTING_MALLOC_DEV( dinfo_magma, magma_int_t, batchCount );
TESTING_MALLOC_DEV( dA_array, magmaFloatComplex*, batchCount );
TESTING_MALLOC_DEV( dipiv_array, magma_int_t*, batchCount );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
// make A diagonally dominant, to not need pivoting
for( int s=0; s < batchCount; ++s ) {
for( int i=0; i < min_mn; ++i ) {
h_A[ i + i*lda + s*lda*N ] = MAGMA_C_MAKE(
MAGMA_C_REAL( h_A[ i + i*lda + s*lda*N ] ) + N,
MAGMA_C_IMAG( h_A[ i + i*lda + s*lda*N ] ));
}
}
columns = N * batchCount;
lapackf77_clacpy( MagmaFullStr, &M, &columns, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, columns, h_R, lda, dA_magma, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_cset_pointer( dA_array, dA_magma, ldda, 0, 0, ldda*N, batchCount, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
info = magma_cgetrf_nopiv_batched( M, N, dA_array, ldda, dinfo_magma, batchCount, opts.queue);
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
// check correctness of results throught "dinfo_magma" and correctness of argument throught "info"
magma_getvector( batchCount, sizeof(magma_int_t), dinfo_magma, 1, cpu_info, 1);
for (int i=0; i < batchCount; i++)
{
if (cpu_info[i] != 0 ) {
printf("magma_cgetrf_batched matrix %d returned internal error %d\n", i, (int)cpu_info[i] );
}
}
if (info != 0) {
printf("magma_cgetrf_batched returned argument error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
for (int i=0; i < batchCount; i++) {
lapackf77_cgetrf(&M, &N, h_A + i*lda*N, &lda, ipiv + i * min_mn, &info);
assert( info == 0 );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f)",
(int) batchCount, (int) M, (int) N, cpu_perf, cpu_time*1000., magma_perf, magma_time*1000., cublas_perf*cublas_enable, cublas_time*1000.*cublas_enable );
}
else {
printf("%10d %5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f (%7.2f)",
(int) batchCount, (int) M, (int) N, magma_perf, magma_time*1000., cublas_perf*cublas_enable, cublas_time*1000.*cublas_enable );
}
if ( opts.check ) {
// initialize ipiv to 1, 2, 3, ...
for (int i=0; i < batchCount; i++)
{
for (int k=0; k < min_mn; k++) {
ipiv[i*min_mn+k] = k+1;
}
}
magma_cgetmatrix( M, N*batchCount, dA_magma, ldda, h_A, lda );
error = 0;
for (int i=0; i < batchCount; i++)
{
float err;
err = get_LU_error( M, N, h_R + i * lda*N, lda, h_A + i * lda*N, ipiv + i * min_mn);
if ( isnan(err) || isinf(err) ) {
error = err;
break;
}
error = max( err, error );
}
bool okay = (error < tol);
status += ! okay;
printf(" %8.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( cpu_info );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( dA_magma );
TESTING_FREE_DEV( dinfo_magma );
TESTING_FREE_DEV( dipiv_magma );
TESTING_FREE_DEV( dipiv_array );
TESTING_FREE_DEV( dA_array );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgesdd (SVD with Divide & Conquer)
Please keep code in testing_cgesdd.cpp and testing_cgesvd.cpp similar.
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
magmaFloatComplex *h_A, *h_R, *U, *VT, *h_work;
magmaFloatComplex dummy[1];
float *S1, *S2;
#ifdef COMPLEX
magma_int_t lrwork=0;
float *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 );
float tol = opts.tolerance * lapackf77_slamch("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_cgesvd_nb( M, N );
// x and y abbreviations used in cgesdd 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_cgesdd( jobz, M, N,
NULL, lda, NULL, NULL, ldu, NULL, ldv, dummy, lwork,
#ifdef COMPLEX
NULL,
#endif
NULL, &info );
lwork = (int) MAGMA_C_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 cgesdd 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, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( VT, magmaFloatComplex, ldv*N ); // N x N (jobz=A) or min(M,N) x N
TESTING_MALLOC_CPU( U, magmaFloatComplex, ldu*N_U ); // M x M (jobz=A) or M x min(M,N)
TESTING_MALLOC_CPU( S1, float, min_mn );
TESTING_MALLOC_CPU( S2, float, min_mn );
TESTING_MALLOC_CPU( iwork, magma_int_t, 8*min_mn );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, 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, float, lrwork );
#endif
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgesdd( 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_cgesdd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
float eps = lapackf77_slamch( "E" );
float 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)
magmaFloatComplex *U2 = NULL;
magmaFloatComplex *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;
}
}
// cbdt01 needs M+N
// cunt01 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) );
}
magmaFloatComplex *h_work_err;
TESTING_MALLOC_CPU( h_work_err, magmaFloatComplex, lwork_err );
// cbdt01 and cunt01 need max(M,N), depending
float *rwork_err;
TESTING_MALLOC_CPU( rwork_err, float, max(M,N) );
if ( U2 != NULL && VT2 != NULL ) {
// since KD=0 (3rd arg), E is not referenced so pass NULL (9th arg)
lapackf77_cbdt01(&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_cunt01("Columns", &M, &N_U, U2, &ldu, h_work_err, &lwork_err,
#ifdef COMPLEX
rwork_err,
#endif
&result[1]);
}
if ( VT2 != NULL ) {
lapackf77_cunt01("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_cgesdd( 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_cgesdd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
float work[1], c_neg_one = -1;
blasf77_saxpy(&min_mn, &c_neg_one, S1, &ione, S2, &ione);
result[4] = lapackf77_slange("f", &min_mn, &ione, S2, &min_mn, work);
result[4] /= lapackf77_slange("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;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sormbr
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, dwork[1];
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t m, n, k, mi, ni, mm, nn, nq, size, info;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t nb, ldc, lda, lwork, lwork_max;
float *C, *R, *A, *work, *tau, *tauq, *taup;
float *d, *e;
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 );
float tol = opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_vect_t vect [] = { MagmaQ, MagmaP };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { MagmaTrans, MagmaNoTrans };
printf(" M N K vect 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 ivect = 0; ivect < 2; ++ivect ) {
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_sgebrd_nb( m );
ldc = m;
// A is nq x k (vect=Q) or k x nq (vect=P)
// where nq=m (left) or nq=n (right)
nq = (side[iside] == MagmaLeft ? m : n );
mm = (vect[ivect] == MagmaQ ? nq : k );
nn = (vect[ivect] == MagmaQ ? k : nq);
lda = mm;
// MBR calls either MQR or MLQ in various ways
if ( vect[ivect] == MagmaQ ) {
if ( nq >= k ) {
gflops = FLOPS_SORMQR( m, n, k, side[iside] ) / 1e9;
}
else {
if ( side[iside] == MagmaLeft ) {
mi = m - 1;
ni = n;
}
else {
mi = m;
ni = n - 1;
}
gflops = FLOPS_SORMQR( mi, ni, nq-1, side[iside] ) / 1e9;
}
}
else {
if ( nq > k ) {
gflops = FLOPS_SORMLQ( m, n, k, side[iside] ) / 1e9;
}
else {
if ( side[iside] == MagmaLeft ) {
mi = m - 1;
ni = n;
}
else {
mi = m;
ni = n - 1;
}
gflops = FLOPS_SORMLQ( mi, ni, nq-1, side[iside] ) / 1e9;
}
}
// workspace for gebrd is (mm + nn)*nb
// workspace for unmbr is m*nb or n*nb, depending on side
lwork_max = max( (mm + nn)*nb, max( m*nb, n*nb ));
TESTING_MALLOC_CPU( C, float, ldc*n );
TESTING_MALLOC_CPU( R, float, ldc*n );
TESTING_MALLOC_CPU( A, float, lda*nn );
TESTING_MALLOC_CPU( work, float, lwork_max );
TESTING_MALLOC_CPU( d, float, min(mm,nn) );
TESTING_MALLOC_CPU( e, float, min(mm,nn) );
TESTING_MALLOC_CPU( tauq, float, min(mm,nn) );
TESTING_MALLOC_CPU( taup, float, min(mm,nn) );
// C is full, m x n
size = ldc*n;
lapackf77_slarnv( &ione, ISEED, &size, C );
lapackf77_slacpy( "Full", &m, &n, C, &ldc, R, &ldc );
size = lda*nn;
lapackf77_slarnv( &ione, ISEED, &size, A );
// compute BRD factorization to get Householder vectors in A, tauq, taup
//lapackf77_sgebrd( &mm, &nn, A, &lda, d, e, tauq, taup, work, &lwork_max, &info );
magma_sgebrd( mm, nn, A, lda, d, e, tauq, taup, work, lwork_max, &info );
if (info != 0)
printf("magma_sgebrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( vect[ivect] == MagmaQ ) {
tau = tauq;
} else {
tau = taup;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_sormbr( lapack_vect_const( vect[ivect] ),
lapack_side_const( side[iside] ),
lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, work, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sormbr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_sormbr( vect[ivect], side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, work, lwork, &info );
if (info != 0)
printf("magma_sormbr (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = (magma_int_t) MAGMA_S_REAL( work[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_sormbr( vect[ivect], side[iside], trans[itran],
m, n, k,
A, lda, tau, R, ldc, work, lwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sormbr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
error = lapackf77_slange( "Fro", &m, &n, C, &ldc, dwork );
size = ldc*n;
blasf77_saxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_slange( "Fro", &m, &n, R, &ldc, dwork ) / error;
printf( "%5d %5d %5d %c %4c %5c %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) m, (int) n, (int) k,
lapacke_vect_const( vect[ivect] ),
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( work );
TESTING_FREE_CPU( d );
TESTING_FREE_CPU( e );
TESTING_FREE_CPU( taup );
TESTING_FREE_CPU( tauq );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}} // end ivect, iside, itran
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e, e1, e2, e3, e4, e5, *work;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, *h_rwork, tmp[1];
magmaFloatComplex_ptr d_A, dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1, ldwork;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// versions 1...4 are valid
if (opts.version < 1 || opts.version > 4) {
printf("Unknown version %d; exiting\n", (int) opts.version );
return -1;
}
float tol = 10. * opts.tolerance * lapackf77_slamch("E");
printf("version %d\n", (int) opts.version );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I-Q'Q||_F / M ||I-Q'Q||_I / M ||A-Q R||_I\n");
printf(" MAGMA / LAPACK MAGMA / LAPACK\n");
printf("==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if (N > 128) {
printf("%5d %5d skipping because cgegqr requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because cgegqr requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
ldwork = N*N;
if (opts.version == 2) {
ldwork = 3*N*N + min_mn + 2;
}
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN(h_rwork, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( work, float, M );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, ldwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
if ( opts.warmup ) {
magma_cgegqr_gpu( 1, M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_cgegqr_gpu( opts.version, M, N, d_A, ldda, dwork, h_rwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
// Regenerate R
// blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_A, &M, &c_zero, h_rwork, &N);
// magma_cprint(N, N, h_work, N);
blasf77_ctrmm("r", "u", "n", "n", &M, &N, &c_one, h_rwork, &N, h_R, &M);
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
e5 = lapackf77_clange("i", &M, &N, h_R, &M, work) /
lapackf77_clange("i", &M, &N, h_A, &lda, work);
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&M, &N, &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_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_cgemm("c", "n", &N, &N, &M, &c_one, h_R, &M, h_R, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e3 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
blasf77_cgemm("c", "n", &N, &N, &M, &c_one, h_A, &M, h_A, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e4 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
if (opts.version != 4)
e = e1;
else
e = e1 / (10.*max(M,N));
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e / %8.2e %8.2e / %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
e1, e2, e3, e4, e5,
(e < tol ? "ok" : "failed"));
status += ! (e < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_rwork );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dsygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time /*cpu_time*/;
double *h_A, *h_R, *h_B, *h_S, *h_work;
double *w1, *w2, vl=0, vu=0;
double result[2] = {0};
magma_int_t *iwork;
magma_int_t N, n2, info, il, iu, m1, m2, nb, lwork, liwork;
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
magma_int_t lrwork;
#endif
//double d_one = 1.;
//double d_ten = 10.;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: itype = %d, jobz = %s, uplo = %s, check = %d, fraction = %6.4f\n",
(int) opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
(int) opts.check, opts.fraction);
printf(" N M GPU Time (sec)\n");
printf("============================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
n2 = N*N;
nb = magma_get_dsytrd_nb(N);
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = 2*N*nb + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = 1 + 6*N*nb + 2* N*N;
#endif
liwork = 3 + 5*N;
if ( opts.fraction == 0 ) {
il = N / 10;
iu = N / 5+il;
}
else {
il = 1;
iu = (int) (opts.fraction*N);
if (iu < 1) iu = 1;
}
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_B, double, n2 );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_PIN( h_S, double, n2 );
TESTING_MALLOC_PIN( h_work, double, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC_PIN( rwork, double, lrwork);
#endif
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlarnv( &ione, ISEED, &n2, h_B );
magma_dmake_hpd( N, h_B, N );
magma_dmake_symmetric( N, h_A, N );
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if(opts.warmup){
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_dsygvdx( opts.itype, opts.jobz, MagmaRangeI, opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_dsygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_dsygvdx( opts.itype, opts.jobz, MagmaRangeI, opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_dsygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | S(with V) - S(w/o V) | / | S |
=================================================================== */
#if defined(PRECISION_d) || defined(PRECISION_s)
double *rwork = h_work + N*N;
#endif
double temp1, temp2;
result[0] = 1.;
result[0] /= lapackf77_dlansy("1", lapack_uplo_const(opts.uplo), &N, h_A, &N, rwork);
result[0] /= lapackf77_dlange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_dscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_dlange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_dscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_dlange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_dscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_dlange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_dsygvdx( opts.itype, MagmaNoVec, MagmaRangeI, opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_dsygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for(int j=0; j < m2; j++) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[1] = temp2 / (((double)m2)*temp1);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %7.2f\n",
(int) N, (int) m1, gpu_time);
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype == 1) {
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
else if (opts.itype == 2) {
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
else if (opts.itype == 3) {
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
printf( "(2) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e %s\n\n", result[1], (result[1] < tolulp ? "ok" : "failed"));
status += ! (result[0] < tol && result[1] < tolulp);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_S );
TESTING_FREE_PIN( h_work );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE_PIN( rwork );
#endif
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing csymmetrize
Code is very similar to testing_ctranspose.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R;
magmaFloatComplex_ptr d_A;
magma_int_t N, size, lda, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N CPU GByte/s (ms) GPU GByte/s (ms) check\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;
size = lda*N;
// load strictly lower triangle, save strictly upper triangle
gbytes = sizeof(magmaFloatComplex) * 1.*N*(N-1) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, size );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, size );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < N; ++i ) {
h_A[i + j*lda] = MAGMA_C_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
gpu_time = magma_sync_wtime( 0 );
//magmablas_csymmetrize( opts.uplo, N-2, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_csymmetrize( opts.uplo, N, d_A, ldda );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using naive in-place algorithm
(LAPACK doesn't implement symmetrize)
=================================================================== */
cpu_time = magma_wtime();
//for( int j = 1; j < N-1; ++j ) { // inset by 1 row & col
// for( int i = 1; i < j; ++i ) {
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < j; ++i ) {
if ( opts.uplo == MagmaLower ) {
h_A[i + j*lda] = MAGMA_C_CNJG( h_A[j + i*lda] );
}
else {
h_A[j + i*lda] = MAGMA_C_CNJG( h_A[i + j*lda] );
}
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
blasf77_caxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &N, &N, h_R, &lda, work);
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
(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" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing stranspose
Code is very similar to testing_ssymmetrize.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, gpu_perf2=0, gpu_time2=0, cpu_perf, cpu_time;
float error, error2, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_B, *h_R;
float *d_A, *d_B;
magma_int_t M, N, size, lda, ldda, ldb, lddb;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("Inplace transpose requires M==N.\n");
printf(" M N CPU GByte/s (ms) GPU GByte/s (ms) check Inplace GB/s (ms) check\n");
printf("====================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = M;
ldda = ((M+31)/32)*32;
ldb = N;
lddb = ((N+31)/32)*32;
// load entire matrix, save entire matrix
gbytes = sizeof(float) * 2.*M*N / 1e9;
TESTING_MALLOC_CPU( h_A, float, lda*N ); // input: M x N
TESTING_MALLOC_CPU( h_B, float, ldb*M ); // output: N x M
TESTING_MALLOC_CPU( h_R, float, ldb*M ); // output: N x M
TESTING_MALLOC_DEV( d_A, float, ldda*N ); // input: M x N
TESTING_MALLOC_DEV( d_B, float, lddb*M ); // output: N x M
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_S_MAKE( i + j/10000., j );
}
}
for( int j = 0; j < M; ++j ) {
for( int i = 0; i < N; ++i ) {
h_B[i + j*ldb] = MAGMA_S_MAKE( i + j/10000., j );
}
}
magma_ssetmatrix( N, M, h_B, ldb, d_B, lddb );
/* =====================================================================
Performs operation using naive out-of-place algorithm
(LAPACK doesn't implement transpose)
=================================================================== */
cpu_time = magma_wtime();
//for( int j = 1; j < N-1; ++j ) { // inset by 1 row & col
// for( int i = 1; i < M-1; ++i ) { // inset by 1 row & col
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_B[j + i*ldb] = h_A[i + j*lda];
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* ====================================================================
Performs operation using MAGMA, out-of-place
=================================================================== */
magma_ssetmatrix( M, N, h_A, lda, d_A, ldda );
magma_ssetmatrix( N, M, h_B, ldb, d_B, lddb );
gpu_time = magma_sync_wtime( 0 );
//magmablas_stranspose( M-2, N-2, d_A+1+ldda, ldda, d_B+1+lddb, lddb ); // inset by 1 row & col
magmablas_stranspose( M, N, d_A, ldda, d_B, lddb );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* ====================================================================
Performs operation using MAGMA, in-place
=================================================================== */
if ( M == N ) {
magma_ssetmatrix( M, N, h_A, lda, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 );
//magmablas_stranspose_inplace( N-2, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_stranspose_inplace( N, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 ) - gpu_time2;
gpu_perf2 = gbytes / gpu_time2;
}
/* =====================================================================
Check the result
=================================================================== */
// check out-of-place transpose (d_B)
size = ldb*M;
magma_sgetmatrix( N, M, d_B, lddb, h_R, ldb );
blasf77_saxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_slange("f", &N, &M, h_R, &ldb, work );
if ( M == N ) {
// also check in-place tranpose (d_A)
magma_sgetmatrix( N, M, d_A, ldda, h_R, ldb );
blasf77_saxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error2 = lapackf77_slange("f", &N, &M, h_R, &ldb, work );
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %6s %7.2f (%7.2f) %s\n",
(int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed"),
gpu_perf2, gpu_time2,
(error2 == 0. ? "ok" : "failed") );
status += ! (error == 0. && error2 == 0.);
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %6s --- ( --- )\n",
(int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cunmqr_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float Cnorm, error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_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, dt_size;
magmaFloatComplex *C, *R, *A, *hwork, *tau;
magmaFloatComplex_ptr dC, dA, dT;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
float tol = opts.tolerance * lapackf77_slamch("E");
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans[] = { Magma_ConjTrans, 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];
ldc = magma_roundup( m, opts.align ); // multiple of 32 by default
// A is m x k (left) or n x k (right)
mm = (side[iside] == MagmaLeft ? m : n);
nb = magma_get_cgeqrf_nb( mm, k );
lda = magma_roundup( mm, opts.align ); // multiple of 32 by default
gflops = FLOPS_CUNMQR( 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;
}
if ( side[iside] == MagmaLeft ) {
// side = left
lwork_max = (m - k + nb)*(n + nb) + n*nb;
dt_size = ( 2*min(m,k) + magma_roundup( max(m,n), 32) )*nb;
}
else {
// side = right
lwork_max = (n - k + nb)*(m + nb) + m*nb;
dt_size = ( 2*min(n,k) + magma_roundup( max(m,n), 32 ) )*nb;
}
// this rounds it up slightly if needed to agree with lwork query below
lwork_max = int( real( magma_cmake_lwork( lwork_max )));
TESTING_MALLOC_CPU( C, magmaFloatComplex, ldc*n );
TESTING_MALLOC_CPU( R, magmaFloatComplex, ldc*n );
TESTING_MALLOC_CPU( A, magmaFloatComplex, lda*k );
TESTING_MALLOC_CPU( hwork, magmaFloatComplex, lwork_max );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, k );
TESTING_MALLOC_DEV( dC, magmaFloatComplex, ldc*n );
TESTING_MALLOC_DEV( dA, magmaFloatComplex, lda*k );
TESTING_MALLOC_DEV( dT, magmaFloatComplex, dt_size );
// C is full, m x n
size = ldc*n;
lapackf77_clarnv( &ione, ISEED, &size, C );
magma_csetmatrix( m, n, C, ldc, dC, ldc );
// A is m x k (left) or n x k (right)
size = lda*k;
lapackf77_clarnv( &ione, ISEED, &size, A );
// compute QR factorization to get Householder vectors in dA, tau, dT
magma_csetmatrix( mm, k, A, lda, dA, lda );
magma_cgeqrf_gpu( mm, k, dA, lda, tau, dT, &info );
magma_cgetmatrix( mm, k, dA, lda, A, lda );
if (info != 0) {
printf("magma_cgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cunmqr( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
&m, &n, &k,
A, &lda, tau, C, &ldc, hwork, &lwork_max, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_cunmqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
// query for workspace size
lwork = -1;
magma_cunmqr_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, hwork, lwork, dT, nb, &info );
if (info != 0) {
printf("magma_cunmqr_gpu (lwork query) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = (magma_int_t) MAGMA_C_REAL( hwork[0] );
if ( lwork < 0 || lwork > lwork_max ) {
printf("Warning: optimal lwork %d > allocated lwork_max %d\n", (int) lwork, (int) lwork_max );
lwork = lwork_max;
}
// cunmqr2 takes a copy of dA in CPU memory
if ( opts.version == 2 ) {
magma_cgetmatrix( mm, k, dA, lda, A, lda );
}
magmablasSetKernelStream( opts.queue );
gpu_time = magma_sync_wtime( opts.queue ); // sync needed for L,N and R,T cases
if ( opts.version == 1 ) {
magma_cunmqr_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, hwork, lwork, dT, nb, &info );
}
else if ( opts.version == 2 ) {
magma_cunmqr2_gpu( side[iside], trans[itran],
m, n, k,
dA, lda, tau, dC, ldc, A, lda, &info );
}
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cunmqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
magma_cgetmatrix( m, n, dC, ldc, R, ldc );
/* =====================================================================
compute relative error |QC_magma - QC_lapack| / |QC_lapack|
=================================================================== */
size = ldc*n;
blasf77_caxpy( &size, &c_neg_one, C, &ione, R, &ione );
Cnorm = lapackf77_clange( "Fro", &m, &n, C, &ldc, work );
error = lapackf77_clange( "Fro", &m, &n, R, &ldc, work ) / (magma_ssqrt(m*n) * Cnorm);
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( hwork );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iside, itran
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cpu_perf, cpu_time;
double magma_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, ldda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ymagma;
magmaDoubleComplex *dA, *dX, *dY, *dwork;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N MAGMA Gflop/s (ms) CPU Gflop/s (ms) MAGMA 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;
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_ZSYMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, sizeY );
blocks = (N + nb - 1) / nb;
ldwork = ldda*blocks;
TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork );
magmablas_zlaset( MagmaFull, ldwork, 1, MAGMA_Z_NAN, MAGMA_Z_NAN, dwork, ldwork );
magmablas_zlaset( MagmaFull, ldda, N, MAGMA_Z_NAN, MAGMA_Z_NAN, dA, ldda );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &sizeA, A );
magma_zmake_hermitian( N, A, lda );
lapackf77_zlarnv( &ione, ISEED, &sizeX, X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, Y );
/* Note: CUBLAS does not implement zsymv */
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_zsetmatrix( N, N, A, lda, dA, ldda );
magma_zsetvector( N, X, incx, dX, incx );
magma_zsetvector( N, Y, incy, dY, incy );
//magma_zprint_gpu( ldda, blocks, dwork, ldda );
magma_time = magma_sync_wtime( 0 );
magmablas_zsymv_work( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy, dwork, ldwork );
// TODO provide option to test non-work interface
//magmablas_zsymv( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( N, dY, incy, Ymagma, incy );
//magma_zprint_gpu( ldda, blocks, dwork, ldda );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zsymv( lapack_uplo_const(opts.uplo), &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_zaxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_zlange( "M", &N, &ione, Ymagma, &N, work ) / N;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N,
magma_perf, 1000.*magma_time,
cpu_perf, 1000.*cpu_time,
magma_error, (magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf
*/
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error;
magmaFloatComplex *h_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t status = 0;
/* Initialize */
magma_queue_t queue[2];
magma_device_t devices[MagmaMaxGPUs];
int num = 0;
magma_err_t err;
magma_init();
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
err = magma_get_devices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
// Create two queues on device opts.device
err = magma_queue_create( devices[opts.device], &queue[0] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( devices[opts.device], &queue[1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |Ax-b|/(N*|A|*|x|)\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |PA-LU|/(N*|A|)\n");
}
printf("=========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGETRF( M, N ) / 1e9;
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_PIN( h_A, magmaFloatComplex, n2 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_cgetrf(&M, &N, h_A, &lda, ipiv, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( M, N, h_A, lda );
gpu_time = magma_wtime();
magma_cgetrf( M, N, h_A, lda, ipiv, &info, queue);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time );
}
if ( opts.check == 2 ) {
error = get_residual( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else if ( opts.check ) {
error = get_LU_error( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_PIN( h_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( queue[0] );
magma_queue_destroy( queue[1] );
magma_finalize();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgetri
*/
int main( int argc, char** argv )
{
TESTING_INIT();
// constants
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double *h_A, *h_Ainv, *h_R, *work;
magmaDouble_ptr d_A, dwork;
magma_int_t N, n2, lda, ldda, info, lwork, ldwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double tmp;
double error, rwork[1];
magma_int_t *ipiv;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% N CPU Gflop/s (sec) GPU Gflop/s (sec) ||I - A*A^{-1}||_1 / (N*cond(A))\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 = magma_roundup( N, opts.align ); // multiple of 32 by default
ldwork = N * magma_get_dgetri_nb( N );
gflops = FLOPS_DGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0) {
printf("lapackf77_dgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = magma_int_t( MAGMA_D_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, double, lwork );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_Ainv, double, n2 );
TESTING_MALLOC_CPU( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( dwork, double, ldwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_dgetrf_gpu( N, N, d_A, ldda, ipiv, &info );
magma_dgetmatrix( N, N, d_A, ldda, h_Ainv, lda );
if (info != 0) {
printf("magma_dgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
// check for exact singularity
//h_Ainv[ 10 + 10*lda ] = MAGMA_D_MAKE( 0.0, 0.0 );
//magma_dsetmatrix( N, N, h_Ainv, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dgetri_gpu( N, d_A, ldda, ipiv, dwork, ldwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_dgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgetri( &N, h_Ainv, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_dgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
printf( "%5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf( "%5d --- ( --- ) %7.2f (%7.2f)",
(int) N, gpu_perf, gpu_time );
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.check ) {
magma_dgetmatrix( N, N, d_A, ldda, h_Ainv, lda );
// compute 1-norm condition number estimate, following LAPACK's zget03
double normA, normAinv, rcond;
normA = lapackf77_dlange( "1", &N, &N, h_A, &lda, rwork );
normAinv = lapackf77_dlange( "1", &N, &N, h_Ainv, &lda, rwork );
if ( normA <= 0 || normAinv <= 0 ) {
rcond = 0;
error = 1 / (tol/opts.tolerance); // == 1/eps
}
else {
rcond = (1 / normA) / normAinv;
// R = I
// R -= A*A^{-1}
// err = ||I - A*A^{-1}|| / ( N ||A||*||A^{-1}|| ) = ||R|| * rcond / N, using 1-norm
lapackf77_dlaset( "full", &N, &N, &c_zero, &c_one, h_R, &lda );
blasf77_dgemm( "no", "no", &N, &N, &N,
&c_neg_one, h_A, &lda,
h_Ainv, &lda,
&c_one, h_R, &lda );
error = lapackf77_dlange( "1", &N, &N, h_R, &lda, rwork );
error = error * rcond / N;
}
bool okay = (error < tol);
status += ! okay;
printf( " %8.2e %s\n",
error, (okay ? "ok" : "failed"));
}
else {
printf( "\n" );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_Ainv );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sswap, sswapblk, slaswp, slaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
float *h_A1, *h_A2;
float *h_R1, *h_R2;
magmaFloat_ptr d_A1, d_A2;
// row-major and column-major performance
real_Double_t row_perf0 = MAGMA_D_NAN, col_perf0 = MAGMA_D_NAN;
real_Double_t row_perf1 = MAGMA_D_NAN, col_perf1 = MAGMA_D_NAN;
real_Double_t row_perf2 = MAGMA_D_NAN, col_perf2 = MAGMA_D_NAN;
real_Double_t row_perf4 = MAGMA_D_NAN;
real_Double_t row_perf5 = MAGMA_D_NAN, col_perf5 = MAGMA_D_NAN;
real_Double_t row_perf6 = MAGMA_D_NAN, col_perf6 = MAGMA_D_NAN;
real_Double_t row_perf7 = MAGMA_D_NAN;
real_Double_t cpu_perf = MAGMA_D_NAN;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magmaInt_ptr d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" %8s sswap sswap sswapblk slaswp slaswp2 slaswpx scopymatrix CPU (all in )\n", g_platform_str );
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj/col-maj row-blk/col-blk slaswp (GByte/s)\n");
printf("=========================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_sgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(float) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, float, lda*N );
TESTING_MALLOC_PIN( h_A2, float, lda*N );
TESTING_MALLOC_PIN( h_R1, float, lda*N );
TESTING_MALLOC_PIN( h_R2, float, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, float, ldda*N );
TESTING_MALLOC_DEV( d_A2, float, ldda*N );
// getrf always makes ipiv[j] >= j+1, where ipiv is one based and j is zero based
// some implementations (e.g., MacOS dlaswp) assume this
for( j=0; j < nb; j++ ) {
ipiv[j] = (rand() % (N-j)) + j + 1;
assert( ipiv[j] >= j+1 );
assert( ipiv[j] <= N );
}
/* =====================================================================
* cublas / clBLAS / Xeon Phi sswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasSswap( opts.handle, N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1 );
#else
magma_sswap( N, d_A1, ldda*j, 1, d_A2, ldda*(ipiv[j]-1), 1, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasSswap( opts.handle, N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
#else
magma_sswap( N, d_A1, j, ldda, d_A2, ipiv[j]-1, ldda, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* sswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_sswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_sswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* sswapblk, blocked version (2 matrices)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_sswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_ssetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_sswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_sgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style slaswp (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style slaswp (1 matrix) - d_ipiv on GPU
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_slaswp2( N, d_A1, ldda, 1, nb, d_ipiv, 1 );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style slaswpx (extended for row- and col-major) (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_sswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_ssetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_slaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
#endif
/* LAPACK swap on CPU for comparison */
time = magma_wtime();
lapackf77_slaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
#ifdef HAVE_CUBLAS
magma_sgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_scopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_scopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf4, ((check & 0x040) != 0 ? '*' : ' '),
row_perf7, ((check & 0x080) != 0 ? '*' : ' '),
row_perf5, ((check & 0x100) != 0 ? '*' : ' '),
col_perf5, ((check & 0x200) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cposv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_perf, cpu_time, gpu_perf, gpu_time;
float error, Rnorm, Anorm, Xnorm, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *h_B, *h_X;
magma_int_t N, lda, ldb, info, sizeA, sizeB;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu = %d, uplo = %s\n", (int) opts.ngpu, lapack_uplo_const(opts.uplo) );
printf(" N NRHS CPU Gflop/s (sec) GPU GFlop/s (sec) ||B - AX|| / N*||A||*||X||\n");
printf("================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = ldb = N;
gflops = ( FLOPS_CPOTRF( N ) + FLOPS_CPOTRS( N, opts.nrhs ) ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*opts.nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*opts.nrhs );
TESTING_MALLOC_CPU( work, float, N );
/* ====================================================================
Initialize the matrix
=================================================================== */
sizeA = lda*N;
sizeB = ldb*opts.nrhs;
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
magma_cmake_hpd( N, h_A, lda );
// copy A to R and B to X; save A and B for residual
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &opts.nrhs, h_B, &ldb, h_X, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cposv( opts.uplo, N, opts.nrhs, h_R, lda, h_X, ldb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Residual
=================================================================== */
Anorm = lapackf77_clange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_clange("I", &N, &opts.nrhs, h_X, &ldb, work);
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &opts.nrhs, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb );
Rnorm = lapackf77_clange("I", &N, &opts.nrhs, h_B, &ldb, work);
error = Rnorm/(N*Anorm*Xnorm);
status += ! (error < tol);
/* ====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cposv( lapack_uplo_const(opts.uplo), &N, &opts.nrhs, h_A, &lda, h_B, &ldb, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cposv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) opts.nrhs, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) opts.nrhs, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( work );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing spotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float *h_A, *h_R;
magma_int_t N, n2, lda, info;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu %d, uplo %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_SPOTRF( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_PIN( h_R, float, n2 );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
magma_smake_hpd( N, h_A, lda );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_spotrf( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_spotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_spotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_spotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_slange("f", &N, &N, h_A, &lda, work);
blasf77_saxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e%s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed") );
status |= ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgels
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float gpu_error, cpu_error, error, Anorm, work[1];
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_A2, *h_B, *h_X, *h_R, *tau, *h_work, tmp[1];
float *d_A, *d_B;
magma_int_t M, N, size, nrhs, lda, ldb, ldda, lddb, min_mn, max_mn, nb, info;
magma_int_t lworkgpu, lhwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
float tol = opts.tolerance * lapackf77_slamch("E");
nrhs = opts.nrhs;
printf(" ||b-Ax|| / (N||A||) ||dx-x||/(N||A||)\n");
printf(" M N NRHS CPU GFlop/s (sec) GPU GFlop/s (sec) CPU GPU \n");
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if ( M < N ) {
printf( "%5d %5d %5d skipping because M < N is not yet supported.\n", (int) M, (int) N, (int) nrhs );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
ldda = ((M+31)/32)*32;
lddb = ((max_mn+31)/32)*32;
nb = magma_get_sgeqrf_nb(M);
gflops = (FLOPS_SGEQRF( M, N ) + FLOPS_SGEQRS( M, N, nrhs )) / 1e9;
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
// query for workspace size
lhwork = -1;
lapackf77_sgels( MagmaNoTransStr, &M, &N, &nrhs,
NULL, &lda, NULL, &ldb, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_S_REAL( tmp[0] );
lhwork = max( lhwork, lworkgpu );
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, lda*N );
TESTING_MALLOC_CPU( h_A2, float, lda*N );
TESTING_MALLOC_CPU( h_B, float, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, float, ldb*nrhs );
TESTING_MALLOC_CPU( h_R, float, ldb*nrhs );
TESTING_MALLOC_CPU( h_work, float, lhwork );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
TESTING_MALLOC_DEV( d_B, float, lddb*nrhs );
/* Initialize the matrices */
size = lda*N;
lapackf77_slarnv( &ione, ISEED, &size, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
size = ldb*nrhs;
lapackf77_slarnv( &ione, ISEED, &size, h_B );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
// make consistent RHS
//size = N*nrhs;
//lapackf77_slarnv( &ione, ISEED, &size, h_X );
//blasf77_sgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
// &c_one, h_A, &lda,
// h_X, &ldb,
// &c_zero, h_B, &ldb );
//lapackf77_slacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix( M, N, h_A, lda, d_A, ldda );
magma_ssetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_sgels_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_work, lworkgpu, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sgels_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute the residual
magma_sgetmatrix( N, nrhs, d_B, lddb, h_X, ldb );
blasf77_sgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb );
Anorm = lapackf77_slange("f", &M, &N, h_A, &lda, work);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_slacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_sgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_work, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_sgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb );
cpu_error = lapackf77_slange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*Anorm);
gpu_error = lapackf77_slange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
// error relative to LAPACK
size = M*nrhs;
blasf77_saxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_slange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e",
(int) M, (int) N, (int) nrhs,
cpu_perf, cpu_time, gpu_perf, gpu_time, cpu_error, gpu_error, error );
if ( M == N ) {
printf( " %s\n", (gpu_error < tol && error < tol ? "ok" : "failed"));
status += ! (gpu_error < tol && error < tol);
}
else {
printf( " %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_A2 );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( h_work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing clanhe
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A;
float *h_work;
magmaFloatComplex_ptr d_A;
magmaFloat_ptr d_work;
magma_int_t i, j, N, n2, lda, ldda;
magma_int_t idist = 3; // normal distribution (otherwise max norm is always ~ 1)
magma_int_t ISEED[4] = {0,0,0,1};
float error, norm_magma, norm_lapack;
magma_int_t status = 0;
magma_int_t lapack_nan_fail = 0;
magma_int_t lapack_inf_fail = 0;
bool mkl_warning = false;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
float tol2;
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper };
magma_norm_t norm[] = { MagmaInfNorm, MagmaOneNorm, MagmaMaxNorm, MagmaFrobeniusNorm };
// Double-Complex inf-norm not supported on Tesla (CUDA arch 1.x)
#if defined(PRECISION_z)
magma_int_t arch = magma_getdevice_arch();
if ( arch < 200 ) {
printf("!!!! NOTE: Double-Complex %s and %s norm are not supported\n"
"!!!! on CUDA architecture %d; requires arch >= 200.\n"
"!!!! It should report \"parameter number 1 had an illegal value\" below.\n\n",
MagmaInfNormStr, MagmaOneNormStr, (int) arch );
for( int inorm = 0; inorm < 2; ++inorm ) {
for( int iuplo = 0; iuplo < 2; ++iuplo ) {
printf( "Testing that magmablas_clanhe( %s, %s, ... ) returns -1 error...\n",
lapack_norm_const( norm[inorm] ),
lapack_uplo_const( uplo[iuplo] ));
norm_magma = magmablas_clanhe( norm[inorm], uplo[iuplo], 1, NULL, 1, NULL, 1 );
if ( norm_magma != -1 ) {
printf( "expected magmablas_clanhe to return -1 error, but got %f\n", norm_magma );
status = 1;
}
}}
printf( "...return values %s\n\n", (status == 0 ? "ok" : "failed") );
}
#endif
#ifdef MAGMA_WITH_MKL
// MKL 11.1 has bug in multi-threaded clanhe; use single thread to work around.
// MKL 11.2 corrects it for inf, one, max norm.
// MKL 11.2 still segfaults for Frobenius norm, which is not tested here
// because MAGMA doesn't implement Frobenius norm yet.
MKLVersion mkl_version;
mkl_get_version( &mkl_version );
magma_int_t la_threads = magma_get_lapack_numthreads();
bool mkl_single_thread = (mkl_version.MajorVersion <= 11 && mkl_version.MinorVersion < 2);
if ( mkl_single_thread ) {
printf( "\nNote: using single thread to work around MKL clanhe bug.\n\n" );
}
#endif
printf("%% N norm uplo CPU GByte/s (ms) GPU GByte/s (ms) error nan inf\n");
printf("%%=================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int inorm = 0; inorm < 3; ++inorm ) { /* < 4 for Frobenius */
for( int iuplo = 0; iuplo < 2; ++iuplo ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = magma_roundup( N, opts.align );
// read upper or lower triangle
gbytes = 0.5*(N+1)*N*sizeof(magmaFloatComplex) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_work, float, N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_work, float, N );
/* Initialize the matrix */
lapackf77_clarnv( &idist, ISEED, &n2, h_A );
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
norm_magma = magmablas_clanhe( norm[inorm], uplo[iuplo], N, d_A, ldda, d_work, N );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gbytes / gpu_time;
if (norm_magma == -1) {
printf( "%5d %4c skipped because %s norm isn't supported\n",
(int) N, lapacke_norm_const( norm[inorm] ), lapack_norm_const( norm[inorm] ));
goto cleanup;
}
else if (norm_magma < 0) {
printf("magmablas_clanhe returned error %f: %s.\n",
norm_magma, magma_strerror( (int) norm_magma ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
#ifdef MAGMA_WITH_MKL
if ( mkl_single_thread ) {
// work around MKL bug in multi-threaded clanhe
magma_set_lapack_numthreads( 1 );
}
#endif
cpu_time = magma_wtime();
norm_lapack = lapackf77_clanhe(
lapack_norm_const( norm[inorm] ),
lapack_uplo_const( uplo[iuplo] ),
&N, h_A, &lda, h_work );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if (norm_lapack < 0) {
printf("lapackf77_clanhe returned error %f: %s.\n",
norm_lapack, magma_strerror( (int) norm_lapack ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = fabs( norm_magma - norm_lapack ) / norm_lapack;
tol2 = tol;
if ( norm[inorm] == MagmaMaxNorm ) {
// max-norm depends on only one element, so for Real precisions,
// MAGMA and LAPACK should exactly agree (tol2 = 0),
// while Complex precisions incur roundoff in cuCabsf.
#ifdef REAL
tol2 = 0;
#endif
}
bool okay; okay = (error <= tol2);
status += ! okay;
mkl_warning |= ! okay;
/* ====================================================================
Check for NAN and INF propagation
=================================================================== */
#define h_A(i_, j_) (h_A + (i_) + (j_)*lda)
#define d_A(i_, j_) (d_A + (i_) + (j_)*ldda)
i = rand() % N;
j = rand() % N;
magma_int_t tmp;
if ( uplo[iuplo] == MagmaLower && i < j ) {
tmp = i;
i = j;
j = tmp;
}
else if ( uplo[iuplo] == MagmaUpper && i > j ) {
tmp = i;
i = j;
j = tmp;
}
*h_A(i,j) = MAGMA_C_NAN;
magma_csetvector( 1, h_A(i,j), 1, d_A(i,j), 1 );
norm_magma = magmablas_clanhe( norm[inorm], uplo[iuplo], N, d_A, ldda, d_work, N );
norm_lapack = lapackf77_clanhe( lapack_norm_const( norm[inorm] ),
lapack_uplo_const( uplo[iuplo] ),
&N, h_A, &lda, h_work );
bool nan_okay; nan_okay = isnan(norm_magma);
bool la_nan_okay; la_nan_okay = isnan(norm_lapack);
lapack_nan_fail += ! la_nan_okay;
status += ! nan_okay;
*h_A(i,j) = MAGMA_C_INF;
magma_csetvector( 1, h_A(i,j), 1, d_A(i,j), 1 );
norm_magma = magmablas_clanhe( norm[inorm], uplo[iuplo], N, d_A, ldda, d_work, N );
norm_lapack = lapackf77_clanhe( lapack_norm_const( norm[inorm] ),
lapack_uplo_const( uplo[iuplo] ),
&N, h_A, &lda, h_work );
bool inf_okay; inf_okay = isinf(norm_magma);
bool la_inf_okay; la_inf_okay = isinf(norm_lapack);
lapack_inf_fail += ! la_inf_okay;
status += ! inf_okay;
#ifdef MAGMA_WITH_MKL
if ( mkl_single_thread ) {
// end single thread to work around MKL bug
magma_set_lapack_numthreads( la_threads );
}
#endif
printf("%5d %4c %4c %7.2f (%7.2f) %7.2f (%7.2f) %#9.3g %-6s %6s%1s %6s%1s\n",
(int) N,
lapacke_norm_const( norm[inorm] ),
lapacke_uplo_const( uplo[iuplo] ),
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
error,
(okay ? "ok" : "failed"),
(nan_okay ? "ok" : "failed"), (la_nan_okay ? " " : "*"),
(inf_okay ? "ok" : "failed"), (la_inf_okay ? " " : "*"));
cleanup:
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_work );
fflush( stdout );
} // end iter
if ( opts.niter > 1 ) {
printf( "\n" );
}
}} // end iuplo, inorm
printf( "\n" );
}
// don't print "failed" here because then run_tests.py thinks MAGMA failed
if ( lapack_nan_fail ) {
printf( "* Warning: LAPACK did not pass NAN propagation test; upgrade to LAPACK version >= 3.4.2 (Sep. 2012)\n" );
}
if ( lapack_inf_fail ) {
printf( "* Warning: LAPACK did not pass INF propagation test\n" );
}
if ( mkl_warning ) {
printf("* MKL (e.g., 11.1) has a bug in clanhe with multiple threads;\n"
" corrected in 11.2 for one, inf, max norms, but still in Frobenius norm.\n"
" Try again with MKL_NUM_THREADS=1.\n" );
}
opts.cleanup();
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;
magmaDouble_ptr 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 );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_DPOTRF( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
magma_dmake_hpd( N, h_A, lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dpotrf_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dpotrf_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, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dlacpy
*/
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_B, *h_R;
magmaDouble_ptr d_A, d_B;
magma_int_t M, N, size, lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("%% uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("%%================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = M;
ldb = lda;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
lddb = ldda;
size = lda*N;
if ( uplo[iuplo] == MagmaLower ) {
// load & save lower trapezoid (with diagonal)
if ( M > N ) {
gbytes = 2. * sizeof(double) * (1.*M*N - 0.5*N*(N-1)) / 1e9;
} else {
gbytes = 2. * sizeof(double) * 0.5*M*(M+1) / 1e9;
}
}
else if ( uplo[iuplo] == MagmaUpper ) {
// load & save upper trapezoid (with diagonal)
if ( N > M ) {
gbytes = 2. * sizeof(double) * (1.*M*N - 0.5*M*(M-1)) / 1e9;
} else {
gbytes = 2. * sizeof(double) * 0.5*N*(N+1) / 1e9;
}
}
else {
// load & save entire matrix
gbytes = 2. * sizeof(double) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, double, size );
TESTING_MALLOC_CPU( h_B, double, size );
TESTING_MALLOC_CPU( h_R, double, size );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( d_B, double, lddb*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 );
h_B[i + j*ldb] = MAGMA_D_MAKE( i - j/10000. + 10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix( M, N, h_A, lda, d_A, ldda, opts.queue );
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
//magmablas_dlacpy( uplo[iuplo], M-2, N-2, d_A+1+ldda, ldda, d_B+1+lddb, lddb, opts.queue ); // inset by 1 row & col
magmablas_dlacpy( uplo[iuplo], M, N, d_A, ldda, d_B, lddb, opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - 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_dlacpy( lapack_uplo_const(uplo[iuplo]), &M2, &N2, h_A+1+lda, &lda, h_B+1+ldb, &ldb );
lapackf77_dlacpy( lapack_uplo_const(uplo[iuplo]), &M, &N, h_A, &lda, h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if ( opts.verbose ) {
printf( "A= " ); magma_dprint( M, N, h_A, lda );
printf( "B= " ); magma_dprint( M, N, h_B, ldb );
printf( "dA=" ); magma_dprint_gpu( M, N, d_A, ldda );
printf( "dB=" ); magma_dprint_gpu( M, N, d_B, lddb );
}
/* =====================================================================
Check the result
=================================================================== */
magma_dgetmatrix( M, N, d_B, lddb, h_R, lda, opts.queue );
blasf77_daxpy(&size, &c_neg_one, h_B, &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_B );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cheevd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
magmaFloatComplex *h_A, *h_R, *h_work, aux_work[1];
float *rwork, *w1, *w2, result[3], eps, aux_rwork[1];
magma_int_t *iwork, aux_iwork[1];
magma_int_t N, n2, info, lwork, lrwork, liwork, lda;
magma_int_t izero = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
eps = lapackf77_slamch( "E" );
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf(" N CPU Time (sec) GPU Time (sec)\n");
printf("=======================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
lda = N;
// query for workspace sizes
magma_cheevd( opts.jobz, opts.uplo,
N, NULL, lda, NULL,
aux_work, -1,
aux_rwork, -1,
aux_iwork, -1,
&info );
lwork = (magma_int_t) MAGMA_C_REAL( aux_work[0] );
lrwork = (magma_int_t) aux_rwork[0];
liwork = aux_iwork[0];
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, N*lda );
TESTING_MALLOC_CPU( w1, float, N );
TESTING_MALLOC_CPU( w2, float, N );
TESTING_MALLOC_CPU( rwork, float, lrwork );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, N*lda );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
magma_cmake_hermitian( N, h_A, N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* warm up run */
if ( opts.warmup ) {
magma_cheevd( opts.jobz, opts.uplo,
N, h_R, lda, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info );
if (info != 0)
printf("magma_cheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cheevd( opts.jobz, opts.uplo,
N, h_R, lda, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_cheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zcds]drvst routine.
A is factored as A = U S U' and the following 3 tests computed:
(1) | A - U S U' | / ( |A| N )
(2) | I - U'U | / ( N )
(3) | S(with U) - S(w/o U) | / | S |
=================================================================== */
float temp1, temp2;
// tau=NULL is unused since itype=1
lapackf77_chet21( &ione, &opts.uplo, &N, &izero,
h_A, &lda,
w1, w1,
h_R, &lda,
h_R, &lda,
NULL, h_work, rwork, &result[0] );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_cheevd( MagmaNoVec, opts.uplo,
N, h_R, lda, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info );
if (info != 0)
printf("magma_cheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for( int j=0; j<N; j++ ) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / (((float)N)*temp1);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cheevd( &opts.jobz, &opts.uplo,
&N, h_A, &lda, w2,
h_work, &lwork,
rwork, &lrwork,
iwork, &liwork,
&info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_cheevd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf("%5d %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time);
}
else {
printf("%5d --- %7.2f\n",
(int) N, gpu_time);
}
/* =====================================================================
Print execution time
=================================================================== */
if ( opts.check ) {
printf("Testing the factorization A = U S U' for correctness:\n");
printf("(1) | A - U S U' | / (|A| N) = %8.2e%s\n", result[0]*eps, (result[0]*eps < tol ? "" : " failed") );
printf("(2) | I - U'U | / N = %8.2e%s\n", result[1]*eps, (result[1]*eps < tol ? "" : " failed") );
printf("(3) | S(w/ U) - S(w/o U) | / |S| = %8.2e%s\n\n", result[2] , (result[2] < tolulp ? "" : " failed") );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( rwork );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dsymmetrize
Code is very similar to testing_dtranspose.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;
magma_int_t N, nb, size, lda, ldda, mstride, nstride, ntile;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
nb = (opts.nb == 0 ? 64 : opts.nb);
mstride = 2*nb;
nstride = 3*nb;
printf("uplo = %s, nb = %d, mstride = %d, nstride = %d\n",
lapack_uplo_const(opts.uplo), (int) nb, (int) mstride, (int) nstride );
printf(" N ntile CPU GByte/s (ms) GPU GByte/s (ms) check\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;
size = lda*N;
if ( N < nb ) {
ntile = 0;
} else {
ntile = min( (N - nb)/mstride + 1,
(N - nb)/nstride + 1 );
}
// load each tile, save each tile
gbytes = sizeof(double) * 2.*nb*nb*ntile / 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 < N; ++i ) {
h_A[i + j*lda] = MAGMA_D_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda );
gpu_time = magma_sync_wtime( 0 );
magmablas_dsymmetrize_tiles( opts.uplo, nb, d_A, ldda, ntile, mstride, nstride );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using naive in-place algorithm
(LAPACK doesn't implement symmetrize)
=================================================================== */
cpu_time = magma_wtime();
for( int tile = 0; tile < ntile; ++tile ) {
int offset = tile*mstride + tile*nstride*lda;
for( int j = 0; j < nb; ++j ) {
for( int i = 0; i < j; ++i ) {
if ( opts.uplo == MagmaLower ) {
h_A[offset + i + j*lda] = MAGMA_D_CNJG( h_A[offset + j + i*lda] );
}
else {
h_A[offset + j + i*lda] = MAGMA_D_CNJG( h_A[offset + i + j*lda] );
}
}
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_dgetmatrix( N, N, d_A, ldda, h_R, lda );
blasf77_daxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("f", &N, &N, h_R, &lda, work);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
(int) N, (int) ntile,
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" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dsygvd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
double *h_A, *h_R, *h_B, *h_S, *h_work;
double *w1, *w2;
magma_int_t *iwork;
magma_int_t N, n2, info, nb, lwork, liwork, lda;
double result[4] = {0};
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double d_zero = 0.;
double d_one = 1.;
double d_neg_one = -1.;
//magma_int_t izero = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf("using: itype = %d, jobz = %s, uplo = %s\n",
(int) opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo));
printf(" N CPU Time (sec) GPU Time(sec)\n");
printf("======================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = N*lda;
nb = magma_get_dsytrd_nb(N);
lwork = 1 + 6*N*nb + 2* N*N;
liwork = 3 + 5*N;
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_B, double, n2 );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( iwork, magma_int_t, liwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_PIN( h_S, double, n2 );
TESTING_MALLOC_PIN( h_work, double, lwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlarnv( &ione, ISEED, &n2, h_B );
magma_dmake_hpd( N, h_B, lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda );
/* warmup */
if ( opts.warmup ) {
magma_dsygvd( opts.itype, opts.jobz, opts.uplo,
N, h_R, lda, h_S, lda, w1,
h_work, lwork,
iwork, liwork,
&info );
if (info != 0)
printf("magma_dsygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dsygvd( opts.itype, opts.jobz, opts.uplo,
N, h_R, lda, h_S, lda, w1,
h_work, lwork,
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_dsygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvd routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | I - V V' B | / ( N ) (itype = 1,2)
| B - V V' | / ( |B| N ) (itype = 3)
(3) | S(with V) - S(w/o V) | / | S |
=================================================================== */
double temp1, temp2;
//double *tau;
if ( opts.itype == 1 || opts.itype == 2 ) {
lapackf77_dlaset( "A", &N, &N, &d_zero, &c_one, h_S, &lda);
blasf77_dgemm("N", "C", &N, &N, &N, &c_one, h_R, &lda, h_R, &lda, &d_zero, h_work, &N);
blasf77_dsymm("R", lapack_uplo_const(opts.uplo), &N, &N, &c_neg_one, h_B, &lda, h_work, &N, &c_one, h_S, &lda);
result[1] = lapackf77_dlange("1", &N, &N, h_S, &lda, h_work) / N;
}
else if ( opts.itype == 3 ) {
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda);
blasf77_dsyrk(lapack_uplo_const(opts.uplo), "N", &N, &N, &d_neg_one, h_R, &lda, &d_one, h_S, &lda);
result[1] = lapackf77_dlansy("1", lapack_uplo_const(opts.uplo), &N, h_S, &lda, h_work) / N
/ lapackf77_dlansy("1", lapack_uplo_const(opts.uplo), &N, h_B, &lda, h_work);
}
result[0] = 1.;
result[0] /= lapackf77_dlansy("1", lapack_uplo_const(opts.uplo), &N, h_A, &lda, h_work);
result[0] /= lapackf77_dlange("1", &N, &N, h_R, &lda, h_work);
if ( opts.itype == 1 ) {
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_A, &lda, h_R, &lda, &d_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_dscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_neg_one, h_B, &lda, h_R, &lda, &c_one, h_work, &N);
result[0] *= lapackf77_dlange("1", &N, &N, h_work, &N, &temp1)/N;
}
else if ( opts.itype == 2 ) {
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_B, &lda, h_R, &lda, &d_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_dscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_A, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_dlange("1", &N, &N, h_R, &lda, &temp1)/N;
}
else if ( opts.itype == 3 ) {
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_A, &lda, h_R, &lda, &d_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_dscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_dsymm("L", lapack_uplo_const(opts.uplo), &N, &N, &c_one, h_B, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_dlange("1", &N, &N, h_R, &lda, &temp1)/N;
}
/*
lapackf77_dsyt21(&ione, lapack_uplo_const(opts.uplo), &N, &izero,
h_A, &lda,
w1, w1,
h_R, &lda,
h_R, &lda,
tau, h_work, rwork, &result[0]);
*/
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda );
magma_dsygvd( opts.itype, MagmaNoVec, opts.uplo,
N, h_R, lda, h_S, lda, w2,
h_work, lwork,
iwork, liwork,
&info );
if (info != 0)
printf("magma_dsygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for(int j=0; j<N; j++) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / temp1;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dsygvd( &opts.itype, lapack_vec_const(opts.jobz), lapack_uplo_const(opts.uplo),
&N, h_A, &lda, h_B, &lda, w2,
h_work, &lwork,
iwork, &liwork,
&info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_dsygvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf("%5d %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time);
}
else {
printf("%5d --- %7.2f\n",
(int) N, gpu_time);
}
/* =====================================================================
Print execution time
=================================================================== */
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if ( opts.itype==1 ) {
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
}
else if ( opts.itype==2 ) {
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
}
else if ( opts.itype==3 ) {
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed") );
}
if ( opts.itype==1 || opts.itype==2 ) {
printf("(2) | I - Z Z' B | / N = %8.2e %s\n", result[1], (result[1] < tol ? "ok" : "failed") );
}
else {
printf("(2) | B - Z Z' | / (|B| N) = %8.2e %s\n", result[1], (result[1] < tol ? "ok" : "failed") );
}
printf( "(3) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e %s\n\n", result[2], (result[2] < tolulp ? "ok" : "failed") );
status += ! (result[0] < tol && result[1] < tol && result[2] < tolulp);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( iwork );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_S );
TESTING_FREE_PIN( h_work );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgemm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
float magma_error, dev_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;
magmaFloatComplex *h_A, *h_B, *h_C, *h_Cmagma, *h_Cdev;
magmaFloatComplex_ptr d_A, d_B, d_C;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.48, 0.38 );
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
#ifdef HAVE_CUBLAS
// for CUDA, we can check MAGMA vs. CUBLAS, without running LAPACK
printf("If running lapack (option --lapack), MAGMA and %s error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to %s result.\n\n",
g_platform_str, g_platform_str );
printf("transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf(" M N K MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
// for others, we need LAPACK for check
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf(" M N K %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
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_CGEMM( 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, magmaFloatComplex, lda*An );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*Bn );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Cmagma, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Cdev, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*An );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
magma_csetmatrix( Am, An, h_A, lda, d_A, ldda );
magma_csetmatrix( Bm, Bn, h_B, ldb, d_B, lddb );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_csetmatrix( M, N, h_C, ldc, d_C, lddc );
magma_time = magma_sync_wtime( NULL );
magmablas_cgemm( 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_cgetmatrix( M, N, d_C, lddc, h_Cmagma, ldc );
#endif
/* =====================================================================
Performs operation using CUBLAS / clBLAS / Xeon Phi MKL
=================================================================== */
magma_csetmatrix( M, N, h_C, ldc, d_C, lddc );
#ifdef HAVE_CUBLAS
dev_time = magma_sync_wtime( NULL );
cublasCgemm( opts.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 );
dev_time = magma_sync_wtime( NULL ) - dev_time;
#else
dev_time = magma_sync_wtime( opts.queue );
magma_cgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, opts.queue );
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
#endif
dev_perf = gflops / dev_time;
magma_cgetmatrix( M, N, d_C, lddc, h_Cdev, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_cgemm( 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 & dev, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_clange( "F", &M, &N, h_C, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cdev, &ione );
dev_error = lapackf77_clange( "F", &M, &N, h_Cdev, &ldc, work ) / Cnorm;
#ifdef HAVE_CUBLAS
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "F", &M, &N, h_Cmagma, &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,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
}
else {
#ifdef HAVE_CUBLAS
// compute relative error for magma, relative to dev (currently only with CUDA)
Cnorm = lapackf77_clange( "F", &M, &N, h_Cdev, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_Cdev, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "F", &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,
dev_perf, 1000.*dev_time,
magma_error,
(magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) --- ( --- ) ---\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time );
#endif
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Cmagma );
TESTING_FREE_CPU( h_Cdev );
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 cpotri
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaFloatComplex *h_A, *h_R;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t N, n2, lda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float Anorm, error, work[1];
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% 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 ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
gflops = FLOPS_CPOTRI( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
magma_cmake_hpd( N, h_A, lda );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
if ( opts.warmup ) {
magma_cpotrf( opts.uplo, N, h_R, lda, &info );
magma_cpotri( opts.uplo, N, h_R, lda, &info );
lapackf77_clacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
}
/* factorize matrix */
magma_cpotrf( opts.uplo, N, h_R, lda, &info );
// check for exact singularity
//h_R[ 10 + 10*lda ] = MAGMA_C_MAKE( 0.0, 0.0 );
gpu_time = magma_wtime();
magma_cpotri( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cpotri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_cpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_cpotri( 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_cpotri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
Anorm = lapackf77_clange("f", &N, &N, h_A, &lda, work);
error = lapackf77_clange("f", &N, &N, h_R, &lda, work) / Anorm;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgemm_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, cublas_error, magma_err, cublas_err, 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;
magma_int_t NN;
magma_int_t batchCount;
magmaFloatComplex *h_A, *h_B, *h_C, *h_Cmagma, *h_Ccublas;
magmaFloatComplex *d_A, *d_B, *d_C;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.48, 0.38 );
magmaFloatComplex **A_array = NULL;
magmaFloatComplex **B_array = NULL;
magmaFloatComplex **C_array = NULL;
magma_queue_t queue = magma_stream;
magma_opts opts;
parse_opts( argc, argv, &opts );
batchCount = opts.batchcount;
cublasHandle_t handle = opts.handle;
//float tol = opts.tolerance * lapackf77_slamch("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"
"transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB));
printf("BatchCount 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_CGEMM( M, N, K ) / 1e9 * batchCount;
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;
NN = N * batchCount;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*An*batchCount;
sizeB = ldb*Bn*batchCount;
sizeC = ldc*N*batchCount;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, sizeA );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, sizeB );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, sizeC );
TESTING_MALLOC_CPU( h_Cmagma, magmaFloatComplex, sizeC );
TESTING_MALLOC_CPU( h_Ccublas, magmaFloatComplex, sizeC );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*An*batchCount );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn*batchCount );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N*batchCount );
magma_malloc((void**)&A_array, batchCount * sizeof(*A_array));
magma_malloc((void**)&B_array, batchCount * sizeof(*B_array));
magma_malloc((void**)&C_array, batchCount * sizeof(*C_array));
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_csetmatrix( Am, An*batchCount, h_A, lda, d_A, ldda );
magma_csetmatrix( Bm, Bn*batchCount, h_B, ldb, d_B, lddb );
magma_csetmatrix( M, N*batchCount, h_C, ldc, d_C, lddc );
cset_pointer(A_array, d_A, ldda, 0, 0, ldda*An, batchCount, queue);
cset_pointer(B_array, d_B, lddb, 0, 0, lddb*Bn, batchCount, queue);
cset_pointer(C_array, d_C, lddc, 0, 0, lddc*N, batchCount, queue);
magma_time = magma_sync_wtime( NULL );
magmablas_cgemm_batched(opts.transA, opts.transB, M, N, K,
alpha, A_array, ldda,
B_array, lddb,
beta, C_array, lddc, batchCount, queue);
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetmatrix( M, N*batchCount, d_C, lddc, h_Cmagma, ldc );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N*batchCount, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasCgemmBatched(handle, cublas_trans_const(opts.transA), cublas_trans_const(opts.transB), M, N, K,
&alpha, (const magmaFloatComplex**) A_array, ldda,
(const magmaFloatComplex**) B_array, lddb,
&beta, C_array, lddc, batchCount );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetmatrix( M, N*batchCount, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
for(int i=0; i<batchCount; i++)
{
blasf77_cgemm(
lapack_trans_const(opts.transA), lapack_trans_const(opts.transB),
&M, &N, &K,
&alpha, h_A + i*lda*An, &lda,
h_B + i*ldb*Bn, &ldb,
&beta, h_C + i*ldc*N, &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|
magma_error = 0.0;
cublas_error = 0.0;
for(int s=0; s<batchCount; s++)
{
magma_int_t C_batchSize = ldc * N;
Cnorm = lapackf77_clange( "M", &M, &N, h_C + s*C_batchSize, &ldc, work );
blasf77_caxpy( &C_batchSize, &c_neg_one, h_C + s*C_batchSize, &ione, h_Cmagma + s*C_batchSize, &ione );
magma_err = lapackf77_clange( "M", &M, &N, h_Cmagma + s*C_batchSize, &ldc, work ) / Cnorm;
if ( isnan(magma_err) || isinf(magma_err) ) {
magma_error = magma_err;
break;
}
magma_error = max(fabs(magma_err), magma_error);
blasf77_caxpy( &C_batchSize, &c_neg_one, h_C + s*C_batchSize, &ione, h_Ccublas + s*C_batchSize, &ione );
cublas_err = lapackf77_clange( "M", &M, &N, h_Ccublas + s*C_batchSize, &ldc, work ) / Cnorm;
if ( isnan(cublas_err) || isinf(cublas_err) ) {
cublas_error = cublas_err;
break;
}
cublas_error = max(fabs(cublas_err), cublas_error);
}
printf("%10d %5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e \n",
(int) batchCount, (int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error);
}
else {
// compute relative error for magma, relative to cublas
Cnorm = lapackf77_clange( "M", &M, &NN, h_Ccublas, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_Ccublas, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "M", &M, &NN, h_Cmagma, &ldc, work ) / Cnorm;
printf("%10d %5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e ---\n",
(int) batchCount, (int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error );
}
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 );
TESTING_FREE_DEV( A_array );
TESTING_FREE_DEV( B_array );
TESTING_FREE_DEV( C_array );
fflush( stdout);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
double magma_error, dev_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ydev, *Ymagma;
magmaDoubleComplex_ptr dA, dX, dY;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("trans = %s\n", lapack_trans_const(opts.transA) );
#ifdef HAVE_CUBLAS
printf(" M N MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
printf(" M N %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = ((M+31)/32)*32;
gflops = FLOPS_ZGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ydev, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, sizeA );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, sizeY );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &sizeA, A );
lapackf77_zlarnv( &ione, ISEED, &sizeX, X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( M, N, A, lda, dA, 0, lda, opts.queue );
magma_zsetvector( Xm, X, incx, dX, 0, incx, opts.queue );
magma_zsetvector( Ym, Y, incy, dY, 0, incy, opts.queue );
#ifdef HAVE_CUBLAS
dev_time = magma_sync_wtime( 0 );
cublasZgemv( opts.handle, cublas_trans_const(opts.transA),
M, N, &alpha, dA, lda, dX, incx, &beta, dY, incy );
dev_time = magma_sync_wtime( 0 ) - dev_time;
#else
dev_time = magma_sync_wtime( opts.queue );
magma_zgemv( opts.transA, M, N,
alpha, dA, 0, lda,
dX, 0, incx,
beta, dY, 0, incy, opts.queue );
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
#endif
dev_perf = gflops / dev_time;
magma_zgetvector( Ym, dY, 0, incy, Ydev, incy, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_zsetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_zgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( Ym, dY, incy, Ymagma, incy );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_zgemv( lapack_trans_const(opts.transA), &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
double Anorm = lapackf77_zlange( "F", &M, &N, A, &lda, work );
double Xnorm = lapackf77_zlange( "F", &Xm, &ione, X, &Xm, work );
blasf77_zaxpy( &Ym, &c_neg_one, Y, &incy, Ydev, &incy );
dev_error = lapackf77_zlange( "F", &Ym, &ione, Ydev, &Ym, work ) / (Anorm * Xnorm);
#ifdef HAVE_CUBLAS
blasf77_zaxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_zlange( "F", &Ym, &ione, Ymagma, &Ym, work ) / (Anorm * Xnorm);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ydev );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
