/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf
*/
int main( magma_int_t argc, char** argv)
{
magma_int_t nquarkthreads=2;
magma_int_t nthreads=2;
magma_int_t num_gpus = 1;
TRACE = 0;
//magma_qr_params mp;
cuDoubleComplex *h_A, *h_R, *h_work, *tau;
double gpu_perf, cpu_perf, flops;
magma_timestr_t start, end;
magma_qr_params *mp = (magma_qr_params*)malloc(sizeof(magma_qr_params));
/* Matrix size */
magma_int_t M=0, N=0, n2;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
cublasStatus status;
magma_int_t i, j, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
mp->nb=-1;
mp->ob=-1;
mp->fb=-1;
mp->ib=32;
magma_int_t loop = argc;
magma_int_t accuracyflag = 1;
char precision;
magma_int_t nc = -1;
magma_int_t ncps = -1;
if (argc != 1)
{
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
else if (strcmp("-F", argv[i])==0)
mp->fb = atoi(argv[++i]);
else if (strcmp("-O", argv[i])==0)
mp->ob = atoi(argv[++i]);
else if (strcmp("-B", argv[i])==0)
mp->nb = atoi(argv[++i]);
else if (strcmp("-b", argv[i])==0)
mp->ib = atoi(argv[++i]);
else if (strcmp("-A", argv[i])==0)
accuracyflag = atoi(argv[++i]);
else if (strcmp("-P", argv[i])==0)
nthreads = atoi(argv[++i]);
else if (strcmp("-Q", argv[i])==0)
nquarkthreads = atoi(argv[++i]);
else if (strcmp("-nc", argv[i])==0)
nc = atoi(argv[++i]);
else if (strcmp("-ncps", argv[i])==0)
ncps = atoi(argv[++i]);
}
if ((M>0 && N>0) || (M==0 && N==0))
{
printf(" testing_zgeqrf-v2 -M %d -N %d\n\n", M, N);
if (M==0 && N==0) {
M = N = size[9];
loop = 1;
}
}
else
{
printf("\nUsage: \n");
printf(" Make sure you set the number of BLAS threads to 1, e.g.,\n");
printf(" > setenv MKL_NUM_THREADS 1\n");
printf(" > testing_zgeqrf-v2 -M %d -N %d -B 128 -T 1\n\n", 1024, 1024);
exit(1);
}
}
else
{
printf("\nUsage: \n");
printf(" Make sure you set the number of BLAS threads to 1, e.g.,\n");
printf(" > setenv MKL_NUM_THREADS 1\n");
printf(" Set number of cores per socket and number of cores.\n");
printf(" > testing_zgeqrf-v2 -M %d -N %d -ncps 6 -nc 12\n\n", 1024, 1024);
printf(" Alternatively, set:\n");
printf(" Q: Number of threads for panel factorization.\n");
printf(" P: Number of threads for trailing matrix update (CPU).\n");
printf(" B: Block size.\n");
printf(" b: Inner block size.\n");
printf(" O: Block size for trailing matrix update (CPU).\n");
printf(" > testing_zgeqrf-v2 -M %d -N %d -Q 4 -P 4 -B 128 -b 32 -O 200\n\n", 10112, 10112);
M = N = size[9];
}
/* Auto tune based on number of cores and number of cores per socket if provided */
if ((nc > 0) && (ncps > 0)) {
precision = 's';
#if (defined(PRECISION_d))
precision = 'd';
#endif
#if (defined(PRECISION_c))
precision = 'c';
#endif
#if (defined(PRECISION_z))
precision = 'z';
#endif
auto_tune('q', precision, nc, ncps, M, N,
&(mp->nb), &(mp->ob), &(mp->ib), &nthreads, &nquarkthreads);
fprintf(stderr,"%d %d %d %d %d\n",mp->nb,mp->ob,mp->ib,nquarkthreads,nthreads);
}
/* Initialize MAGMA hardware context, seeting how many CPU cores
and how many GPUs to be used in the consequent computations */
mp->sync0 = 0;
magma_context *context;
context = magma_init((void*)(mp),cpu_thread, nthreads, nquarkthreads, num_gpus, argc, argv);
context->params = (void *)(mp);
mp->sync1 = (volatile magma_int_t *) malloc (sizeof(int)*nthreads);
for (i = 0; i < nthreads; i++)
mp->sync1[i] = 0;
n2 = M * N;
magma_int_t min_mn = min(M, N);
magma_int_t nb = magma_get_zgeqrf_nb(min_mn);
magma_int_t lwork = N*nb;
/* Allocate host memory for the matrix */
TESTING_MALLOC ( h_A , cuDoubleComplex, n2 );
TESTING_MALLOC ( tau , cuDoubleComplex, min_mn);
TESTING_HOSTALLOC( h_R , cuDoubleComplex, n2 );
TESTING_HOSTALLOC(h_work, cuDoubleComplex, lwork );
printf("\n\n");
printf(" M N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("==========================================================\n");
for(i=0; i<10; i++){
if (loop==1){
M = N = min_mn = size[i];
n2 = M*N;
}
flops = FLOPS( (double)M, (double)N ) / 1000000;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &M, h_R, &M );
//magma_zgeqrf(M, N, h_R, M, tau, h_work, lwork, &info);
for(j=0; j<n2; j++)
h_R[j] = h_A[j];
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_qr_init(mp, M, N, h_R, nthreads);
start = get_current_time();
magma_zgeqrf3(context, M, N, h_R, M, tau, h_work, lwork, &info);
end = get_current_time();
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
if (accuracyflag == 1)
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, h_work, &lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of zgeqrf had an illegal value.\n", -info);
cpu_perf = 4.*M*N*min_mn/(3.*1000000*GetTimerValue(start,end));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
double work[1], matnorm = 1.;
cuDoubleComplex mone = MAGMA_Z_NEG_ONE;
magma_int_t one = 1;
if (accuracyflag == 1){
matnorm = lapackf77_zlange("f", &M, &N, h_A, &M, work);
blasf77_zaxpy(&n2, &mone, h_A, &one, h_R, &one);
}
if (accuracyflag == 1){
printf("%5d %5d %6.2f %6.2f %e\n",
M, N, cpu_perf, gpu_perf,
lapackf77_zlange("f", &M, &N, h_R, &M, work) / matnorm);
} else {
printf("%5d %5d %6.2f \n",
M, N, gpu_perf);
}
if (loop != 1)
break;
}
/* Memory clean up */
TESTING_FREE ( h_A );
TESTING_FREE ( tau );
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R );
/* Shut down the MAGMA context */
magma_finalize(context);
}
int main(int argc, char **argv)
{
TESTING_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 zpotf2_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magma_int_t N, n2, lda, ldda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_ZPOTRF( N ) / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zpotf2_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
error = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., error );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
// --------------------
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_time=0, cpu_perf=0, gpu_time, gpu_perf, mgpu_time, mgpu_perf, cuda_time, cuda_perf;
double error=0, error2=0, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t n_local[MagmaMaxGPUs];
magma_int_t N, Noffset, lda, ldda, blocks, lhwork, ldwork, matsize, vecsize;
magma_int_t incx = 1;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ylapack, *Ycublas, *Ymagma, *Ymagma1, *hwork;
magmaDoubleComplex_ptr dA, dX, dY;
magmaDoubleComplex_ptr d_lA[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magma_device_t dev;
magma_queue_t queues[MagmaMaxGPUs];
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
magma_int_t nb = 64; // required by magmablas_zhemv_mgpu implementation
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_create( &queues[dev] );
}
// currently, tests all offsets in the offsets array;
// comment out loop below to test a specific offset.
magma_int_t offset = opts.offset;
magma_int_t offsets[] = { 0, 1, 31, 32, 33, 63, 64, 65, 100, 200 };
magma_int_t noffsets = sizeof(offsets) / sizeof(*offsets);
printf("uplo = %s, ngpu %d, block size = %d, offset %d\n",
lapack_uplo_const(opts.uplo), (int) opts.ngpu, (int) nb, (int) offset );
printf( " BLAS CUBLAS MAGMA 1 GPU MAGMA MGPU Error rel Error rel\n"
" N offset Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) to CUBLAS to LAPACK\n"
"===================================================================================================================\n" );
for( int itest = 0; itest < opts.ntest; ++itest ) {
// comment out these two lines & end of loop to test a specific offset
for( int ioffset=0; ioffset < noffsets; ioffset += 1 ) {
offset = offsets[ioffset];
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
Noffset = N + offset;
lda = Noffset;
ldda = ((Noffset+31)/32)*32;
matsize = Noffset*ldda;
vecsize = (Noffset-1)*incx + 1;
gflops = FLOPS_ZHEMV( N ) / 1e9;
blocks = (N + (offset % nb) - 1)/nb + 1;
lhwork = N*opts.ngpu;
ldwork = ldda*(blocks + 1);
TESTING_MALLOC_CPU( A, magmaDoubleComplex, matsize );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ycublas, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma1, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ylapack, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( X, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( hwork, magmaDoubleComplex, lhwork );
magma_setdevice( opts.device );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, matsize );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, vecsize );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, vecsize );
// TODO make magma_zmalloc_bcyclic helper function?
for( dev=0; dev < opts.ngpu; dev++ ) {
n_local[dev] = ((Noffset/nb)/opts.ngpu)*nb;
if (dev < (Noffset/nb) % opts.ngpu)
n_local[dev] += nb;
else if (dev == (Noffset/nb) % opts.ngpu)
n_local[dev] += Noffset % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaDoubleComplex, ldda*n_local[dev] );
TESTING_MALLOC_DEV( dwork[dev], magmaDoubleComplex, ldwork );
}
//////////////////////////////////////////////////////////////////////////
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &matsize, A );
magma_zmake_hermitian( Noffset, A, lda );
lapackf77_zlarnv( &ione, ISEED, &vecsize, X );
lapackf77_zlarnv( &ione, ISEED, &vecsize, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_setdevice( opts.device );
magma_zsetmatrix( Noffset, Noffset, A, lda, dA, ldda );
magma_zsetvector( Noffset, X, incx, dX, incx );
magma_zsetvector( Noffset, Y, incx, dY, incx );
cuda_time = magma_sync_wtime(0);
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo), N,
&alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
&beta, dY + offset, incx );
cuda_time = magma_sync_wtime(0) - cuda_time;
cuda_perf = gflops / cuda_time;
magma_zgetvector( Noffset, dY, incx, Ycublas, incx );
/* =====================================================================
Performs operation using MAGMABLAS (1 GPU)
=================================================================== */
magma_setdevice( opts.device );
magma_zsetvector( Noffset, Y, incx, dY, incx );
gpu_time = magma_sync_wtime( opts.queue );
magmablas_zhemv_work( opts.uplo, N,
alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
beta, dY + offset, incx, dwork[ opts.device ], ldwork,
opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_zgetvector( Noffset, dY, incx, Ymagma1, incx );
/* =====================================================================
Performs operation using MAGMABLAS (multi-GPU)
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( Noffset, Noffset, A, lda, d_lA, ldda, opts.ngpu, nb );
blasf77_zcopy( &Noffset, Y, &incx, Ymagma, &incx );
// workspaces do NOT need to be zero -- set to NAN to prove
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magmablas_zlaset( MagmaFull, ldwork, 1, MAGMA_Z_NAN, MAGMA_Z_NAN, dwork[dev], ldwork );
}
lapackf77_zlaset( "Full", &lhwork, &ione, &MAGMA_Z_NAN, &MAGMA_Z_NAN, hwork, &lhwork );
mgpu_time = magma_sync_wtime(0);
magma_int_t info;
info = magmablas_zhemv_mgpu(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if ( info != 0 )
printf("magmablas_zhemv_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
info = magmablas_zhemv_mgpu_sync(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if ( info != 0 )
printf("magmablas_zhemv_sync returned error %d: %s.\n",
(int) info, magma_strerror( info ));
mgpu_time = magma_sync_wtime(0) - mgpu_time;
mgpu_perf = gflops / mgpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
blasf77_zcopy( &Noffset, Y, &incx, Ylapack, &incx );
cpu_time = magma_wtime();
blasf77_zhemv( lapack_uplo_const(opts.uplo), &N,
&alpha, A + offset + offset*lda, &lda,
X + offset, &incx,
&beta, Ylapack + offset, &incx );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Compute the Difference LAPACK vs. Magma
=================================================================== */
error2 = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ylapack, &incx );
error2 = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work ) / error2;
}
/* =====================================================================
Compute the Difference Cublas vs. Magma
=================================================================== */
error = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ycublas, &incx );
error = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work ) / error;
bool okay = (error < tol && error2 < tol);
status += ! okay;
if ( opts.lapack ) {
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) N, (int) offset,
cpu_perf, cpu_time*1000.,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, error2, (okay ? "ok" : "failed") );
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e --- %s\n",
(int) N, (int) offset,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, (okay ? "ok" : "failed") );
}
/* Free Memory */
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_CPU( Ymagma1 );
TESTING_FREE_CPU( Ylapack );
TESTING_FREE_PIN( X );
TESTING_FREE_PIN( hwork );
magma_setdevice( opts.device );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
for( dev=0; dev < opts.ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
// comment out these two lines line & top of loop test a specific offset
} // end for ioffset
printf( "\n" );
}
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_destroy( queues[dev] );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double magma_error, cublas_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *piv;
magma_err_t err;
magmaDoubleComplex *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2, *LU, *LUT;
magmaDoubleComplex *d_A, *d_B;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"side = %c, uplo = %c, transA = %c, diag = %c \n", opts.side, opts.uplo, opts.transA, opts.diag );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("==================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
gflops = FLOPS_ZTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( LU, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( LUT, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_B1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_X1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_X2, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_Bcublas, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_Bmagma, magmaDoubleComplex, ldb*N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, LU );
err = magma_malloc_cpu( (void**) &piv, Ak*sizeof(magma_int_t) ); assert( err == 0 );
lapackf77_zgetrf( &Ak, &Ak, LU, &lda, piv, &info );
int i, j;
for(i=0;i<Ak;i++){
for(j=0;j<Ak;j++){
LUT[j+i*lda] = LU[i+j*lda];
}
}
lapackf77_zlacpy(MagmaUpperStr, &Ak, &Ak, LUT, &lda, LU, &lda);
if(opts.uplo == MagmaLower){
lapackf77_zlacpy(MagmaLowerStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}else{
lapackf77_zlacpy(MagmaUpperStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(magmaDoubleComplex));
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_zsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_ztrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasZtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ztrsm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(magmaDoubleComplex));
magmaDoubleComplex alpha2 = MAGMA_Z_DIV( c_one, alpha );
blasf77_ztrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
double norm1 = lapackf77_zlange( "M", &M, &N, h_X1, &ldb, work );
double normx = lapackf77_zlange( "M", &M, &N, h_Bmagma, &ldb, work );
double normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(magmaDoubleComplex));
blasf77_ztrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_zlange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_zlange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error );
}
TESTING_FREE( h_A );
TESTING_FREE( LU );
TESTING_FREE( LUT );
TESTING_FREE( h_B );
TESTING_FREE( h_Bcublas );
TESTING_FREE( h_Bmagma );
TESTING_FREE( h_B1 );
TESTING_FREE( h_X1 );
TESTING_FREE( h_X2 );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_zlahr2(magma_int_t n, magma_int_t k, magma_int_t nb,
cuDoubleComplex *da, cuDoubleComplex *dv,
cuDoubleComplex *a, magma_int_t lda,
cuDoubleComplex *tau, cuDoubleComplex *t, magma_int_t ldt,
cuDoubleComplex *y, magma_int_t ldy)
{
/* -- MAGMA auxiliary routine (version 1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
ZLAHR2 reduces the first NB columns of a complex general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
This is an auxiliary routine called by ZGEHRD.
Arguments
=========
N (input) INTEGER
The order of the matrix A.
K (input) INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
NB (input) INTEGER
The number of columns to be reduced.
DA (input/output) COMPLEX_16 array on the GPU, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
DV (output) COMPLEX_16 array on the GPU, dimension (N, NB)
On exit this contains the Householder vectors of the transformation.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX_16 array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) COMPLEX_16 array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) COMPLEX_16 array, dimension (LDY,NB)
The n-by-nb matrix Y.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
Further Details
===============
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
===================================================================== */
cuDoubleComplex c_zero = MAGMA_Z_ZERO;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ldda = lda;
magma_int_t c__1 = 1;
magma_int_t a_dim1, a_offset, t_dim1, t_offset, y_dim1, y_offset, i__2, i__3;
cuDoubleComplex d__1;
magma_int_t i__;
cuDoubleComplex ei;
--tau;
a_dim1 = lda;
a_offset = 1 + a_dim1;
a -= a_offset;
t_dim1 = ldt;
t_offset = 1 + t_dim1;
t -= t_offset;
y_dim1 = ldy;
y_offset = 1 + y_dim1;
y -= y_offset;
/* Function Body */
if (n <= 1)
return 0;
for (i__ = 1; i__ <= nb; ++i__) {
if (i__ > 1) {
/* Update A(K+1:N,I); Update I-th column of A - Y * V' */
i__2 = n - k + 1;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
blasf77_zcopy(&i__3, &a[k+i__-1+a_dim1], &lda, &t[nb*t_dim1+1], &c__1);
blasf77_ztrmv("u","n","n",&i__3,&t[t_offset], &ldt, &t[nb*t_dim1+1], &c__1);
blasf77_zgemv("NO TRANSPOSE", &i__2, &i__3, &c_neg_one, &y[k + y_dim1],
&ldy, &t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__*a_dim1],&c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
/* Apply I - V * T' * V' to this column (call it b) from the
left, using the last column of T as workspace
Let V = ( V1 ) and b = ( b1 ) (first I-1 rows)
( V2 ) ( b2 )
where V1 is unit lower triangular
w := V1' * b1 */
i__2 = i__ - 1;
blasf77_zcopy(&i__2, &a[k+1+i__*a_dim1], &c__1, &t[nb*t_dim1+1], &c__1);
blasf77_ztrmv("Lower", MagmaConjTransStr, "UNIT", &i__2,
&a[k + 1 + a_dim1], &lda, &t[nb * t_dim1 + 1], &c__1);
/* w := w + V2'*b2 */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_one, &t[nb*t_dim1+1], &c__1);
/* w := T'*w */
i__2 = i__ - 1;
blasf77_ztrmv("U", MagmaConjTransStr, "N", &i__2, &t[t_offset], &ldt,
&t[nb*t_dim1+1], &c__1);
/* b2 := b2 - V2*w */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv("N", &i__2, &i__3, &c_neg_one, &a[k + i__ + a_dim1], &lda,
&t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__+i__*a_dim1], &c__1);
/* b1 := b1 - V1*w */
i__2 = i__ - 1;
blasf77_ztrmv("L","N","U",&i__2,&a[k+1+a_dim1],&lda,&t[nb*t_dim1+1],&c__1);
blasf77_zaxpy(&i__2, &c_neg_one, &t[nb * t_dim1 + 1], &c__1,
&a[k + 1 + i__ * a_dim1], &c__1);
a[k + i__ - 1 + (i__ - 1) * a_dim1] = ei;
}
/* Generate the elementary reflector H(I) to annihilate A(K+I+1:N,I) */
i__2 = n - k - i__ + 1;
i__3 = k + i__ + 1;
lapackf77_zlarfg(&i__2, &a[k + i__ + i__ * a_dim1],
&a[min(i__3,n) + i__ * a_dim1], &c__1, &tau[i__]);
ei = a[k + i__ + i__ * a_dim1];
a[k + i__ + i__ * a_dim1] = c_one;
/* Compute Y(K+1:N,I) */
i__2 = n - k;
i__3 = n - k - i__ + 1;
magma_zsetvector( i__3,
&a[k + i__ + i__*a_dim1], 1,
dv+(i__-1)*(ldda+1), 1 );
magma_zgemv(MagmaNoTrans, i__2+1, i__3, c_one,
da -1 + k + i__ * ldda, ldda,
dv+(i__-1)*(ldda+1), c__1, c_zero,
da-1 + k + (i__-1)*ldda, c__1);
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_zero, &t[i__*t_dim1+1], &c__1);
/* Compute T(1:I,I) */
i__2 = i__ - 1;
d__1 = MAGMA_Z_NEGATE( tau[i__] );
blasf77_zscal(&i__2, &d__1, &t[i__ * t_dim1 + 1], &c__1);
blasf77_ztrmv("U","N","N", &i__2, &t[t_offset], &ldt, &t[i__*t_dim1+1], &c__1);
t[i__ + i__ * t_dim1] = tau[i__];
magma_zgetvector( n - k + 1,
da-1+ k+(i__-1)*ldda, 1,
y+ k + i__*y_dim1, 1 );
}
a[k + nb + nb * a_dim1] = ei;
return 0;
} /* magma_zlahr2 */
/**
Purpose
-------
ZLATRD reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in]
nb0 INTEGER
The block size used for the matrix distribution.
nb and nb0 can be different for the final step of zhetrd.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX_16 array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
@param
dA
@param[in]
ldda
@param[in]
offset
@param
dW
@param[in]
lddw
@param
hwork
@param[in]
lhwork
@param
dwork
@param[in]
ldwork
@param[in]
queues magma_queue_t array of dimension (ngpu).
queues[dev] is an execution queue on GPU dev.
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo,
magma_int_t n, magma_int_t nb, magma_int_t nb0,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex_ptr dA[], magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex_ptr dW[], magma_int_t lddw,
magmaDoubleComplex *hwork, magma_int_t lhwork,
magmaDoubleComplex_ptr dwork[], magma_int_t ldwork,
magma_queue_t queues[] )
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(dev, i, j) (dA[(dev)] + ((j)+loffset)*ldda + (i) + offset)
#define dW(dev, i, j) (dW[(dev)] + (j) *lddw + (i))
#define dW1(dev, i, j) (dW[(dev)] + ((j)+nb) *lddw + (i))
/* Constants */
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magma_int_t ione = 1;
/* Local variables */
magmaDoubleComplex alpha, value;
magma_int_t dev;
magma_int_t i, n_i, n_i_1, ip1, iw;
// TODO check arguments
magma_int_t info = 0;
if (n <= 0) {
return info;
}
// TODO allocate f in zhetrd and pass into zlatrd. (e.g., expand hwork a bit)
magmaDoubleComplex *f;
magma_zmalloc_cpu( &f, n );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
return info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
ip1 = i + 1;
n_i_1 = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
magmaDoubleComplex wii = -conj( *W(i, iw+1) );
blasf77_zaxpy( &ip1, &wii, A(0, i+1), &ione, A(0, i), &ione );
wii = -conj( *A(i, i+1) );
blasf77_zaxpy( &ip1, &wii, W(0, iw+1), &ione, A(0, i), &ione );
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_ONE;
// TODO Previously, this set dx2[dev] = dW1(dev, 0, iw); and used dx2 in zhemv.
// TODO Now zhemv handles broadcasting x to the GPUs, but data in dW1 is
// TODO apparently still used in zhetrd_mgpu / zher2k_mgpu.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, A(0,i), 1, dW1(dev, 0, iw), 1, queues[dev] );
}
magmablas_zhemv_mgpu(
MagmaUpper, i, c_one, dA, ldda, 0,
A(0,i), 1, c_zero, W(0, iw), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
if (i < n-1) {
blasf77_zgemv( MagmaConjTransStr, &i, &n_i_1, &c_one,
W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero,
W(i+1, iw), &ione );
}
/* overlap update */
if ( i < n-1 && i-1 >= n - nb ) {
/* Update A(1:i,i) */
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, W(i-1, iw+1), &ldw );
#endif
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
A(0, i+1), &lda,
W(i-1, iw+1), &ldw, &c_one,
A(0, i-1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, W(i-1, iw+1), &ldw );
lapackf77_zlacgv( &n_i_1, A(i-1, i +1), &lda );
#endif
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
W(0, iw+1), &ldw,
A(i-1, i+1), &lda, &c_one,
A(0, i-1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, A(i-1, i+1), &lda );
#endif
}
// synchronize to get zhemv result W(0, iw)
magmablas_zhemv_mgpu_sync(
MagmaUpper, i, c_one, dA, ldda, 0,
A(0,i), 1, c_zero, W(0, iw), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
if (i < n-1) {
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one,
W(0, iw), &ione );
blasf77_zgemv( MagmaConjTransStr, &i, &n_i_1, &c_one,
A(0, i+1), &lda,
A(0, i), &ione, &c_zero,
W(i+1, iw), &ione );
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one,
W(0, iw), &ione );
}
blasf77_zscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy( &i, &alpha, A(0, i), &ione, W(0, iw), &ione );
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, W(0,iw), 1, dW(dev, 0, iw), 1, queues[dev] );
}
}
}
} else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
n_i = n - i;
//idw = ((offset+i)/nb)%ngpu;
if ( i > 0 ) {
trace_cpu_start( 0, "gemv", "gemv" );
magmaDoubleComplex wii = -conj( *W(i, i-1) );
blasf77_zaxpy( &n_i, &wii, A(i, i-1), &ione, A(i, i), &ione );
wii = -conj( *A(i, i-1) );
blasf77_zaxpy( &n_i, &wii, W(i, i-1), &ione, A(i, i), &ione );
}
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
n_i_1 = n - i - 1;
trace_cpu_start( 0, "larfg", "larfg" );
alpha = *A(i+1, i);
lapackf77_zlarfg( &n_i_1, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_ONE;
trace_cpu_end( 0 );
/* Compute W(i+1:n,i) */
// TODO Previously, this set dx2[id] = dW1(id, 0, i)-offset; and used dx2 in zhemv.
// TODO Now zhemv handles broadcasting x to the GPUs, but data in dW1 is
// TODO apparently still used in zhetrd_mgpu / zher2k_mgpu.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, A(0,i), 1, dW1(dev, 0, i), 1, queues[dev] );
}
magmablas_zhemv_mgpu(
MagmaLower, n_i_1, c_one, dA, ldda, offset+i+1,
A(i+1, i), 1, c_zero, W(i+1, i), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
trace_cpu_start( 0, "gemv", "gemv" );
blasf77_zgemv( MagmaConjTransStr, &n_i_1, &i, &c_one,
W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero,
W(0, i), &ione );
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
A(i+1, 0), &lda,
W(0, i), &ione, &c_zero,
f, &ione );
blasf77_zgemv( MagmaConjTransStr, &n_i_1, &i, &c_one,
A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero,
W(0, i), &ione );
trace_cpu_end( 0 );
/* overlap update */
if ( i > 0 && i+1 < n ) {
trace_cpu_start( 0, "gemv", "gemv" );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i+1, 0), &ldw );
#endif
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
A(i+1, 0), &lda,
W(i+1, 0), &ldw, &c_one,
A(i+1, i+1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i+1, 0), &ldw );
lapackf77_zlacgv( &i, A(i+1, 0), &lda );
#endif
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
W(i+1, 0), &ldw,
A(i+1, 0), &lda, &c_one,
A(i+1, i+1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, A(i+1, 0), &lda );
#endif
trace_cpu_end( 0 );
}
// synchronize to get zhemv result W(i+1, i)
magmablas_zhemv_mgpu_sync(
MagmaLower, n_i_1, c_one, dA, ldda, offset+i+1,
A(i+1, i), 1, c_zero, W(i+1, i), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
trace_cpu_start( 0, "axpy", "axpy" );
if (i != 0) {
blasf77_zaxpy( &n_i_1, &c_one, f, &ione, W(i+1, i), &ione );
}
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
W(i+1, 0), &ldw,
W(0, i), &ione, &c_one,
W(i+1, i), &ione );
blasf77_zscal( &n_i_1, &tau[i], W(i+1,i), &ione );
value = magma_cblas_zdotc( n_i_1, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy( &n_i_1, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
trace_cpu_end( 0 );
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, W(0,i), 1, dW(dev, 0, i), 1, queues[dev] );
}
}
}
}
magma_free_cpu( f );
magma_setdevice( orig_dev );
return info;
} /* magma_zlatrd_mgpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgemm_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cpu_perf, cpu_time;
double magma_error, magma_err, Ynorm, work[1];
magma_int_t M, N, Xm, Ym, lda, ldda;
magma_int_t sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_int_t batchCount;
magmaDoubleComplex *h_A, *h_X, *h_Y, *h_Ymagma;
magmaDoubleComplex *d_A, *d_X, *d_Y;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.48, 0.38 );
magmaDoubleComplex **A_array = NULL;
magmaDoubleComplex **X_array = NULL;
magmaDoubleComplex **Y_array = NULL;
magma_opts opts;
parse_opts( argc, argv, &opts );
batchCount = opts.batchcount;
opts.lapack |= opts.check;
//double tol = opts.tolerance * lapackf77_dlamch("E");
printf("trans = %s\n", lapack_trans_const(opts.transA) );
printf("BatchCount M 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 ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = ((M+31)/32)*32;
gflops = FLOPS_ZGEMV( M, N ) / 1e9 * batchCount;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N*batchCount;
sizeX = incx*Xm*batchCount;
sizeY = incy*Ym*batchCount;
ldda = ((lda+31)/32)*32;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( h_X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( h_Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( h_Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N*batchCount );
TESTING_MALLOC_DEV( d_X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( d_Y, magmaDoubleComplex, sizeY );
magma_malloc((void**)&A_array, batchCount * sizeof(*A_array));
magma_malloc((void**)&X_array, batchCount * sizeof(*X_array));
magma_malloc((void**)&Y_array, batchCount * sizeof(*Y_array));
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &sizeX, h_X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, h_Y );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_zsetmatrix( M, N*batchCount, h_A, lda, d_A, ldda );
magma_zsetvector( Xm*batchCount, h_X, incx, d_X, incx );
magma_zsetvector( Ym*batchCount, h_Y, incy, d_Y, incy );
zset_pointer(A_array, d_A, ldda, 0, 0, ldda*N, batchCount, magma_stream);
zset_pointer(X_array, d_X, 1, 0, 0, incx*Xm, batchCount, magma_stream);
zset_pointer(Y_array, d_Y, 1, 0, 0, incy*Ym, batchCount, magma_stream);
magma_time = magma_sync_wtime( NULL );
magmablas_zgemv_batched(opts.transA, M, N,
alpha, A_array, ldda,
X_array, incx,
beta, Y_array, incy, batchCount, magma_stream);
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( Ym*batchCount, d_Y, incy, h_Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
for(int i=0; i<batchCount; i++)
{
blasf77_zgemv(
lapack_trans_const(opts.transA),
&M, &N,
&alpha, h_A + i*lda*N, &lda,
h_X + i*Xm, &incx,
&beta, h_Y + i*Ym, &incy );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
magma_error = 0.0;
for(int s=0; s<batchCount; s++)
{
Ynorm = lapackf77_zlange( "M", &M, &ione, h_Y + s*Ym, &incy, work );
blasf77_zaxpy( &Ym, &c_neg_one, h_Y + s*Ym, &ione, h_Ymagma + s*Ym, &ione );
magma_err = lapackf77_zlange( "M", &M, &ione, h_Ymagma + s*Ym, &incy, work ) / Ynorm;
if ( isnan(magma_err) || isinf(magma_err) ) {
magma_error = magma_err;
break;
}
magma_error = max(fabs(magma_err), magma_error);
}
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e \n",
(int) batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time,
cpu_perf, 1000.*cpu_time,
magma_error);
}
else {
printf("%10d %5d %5d %7.2f (%7.2f) --- ( --- ) ---\n",
(int) batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_Y );
TESTING_FREE_CPU( h_Ymagma );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_X );
TESTING_FREE_DEV( d_Y );
TESTING_FREE_DEV( A_array );
TESTING_FREE_DEV( X_array );
TESTING_FREE_DEV( Y_array );
fflush( stdout);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeev
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
double *h_A, *h_R, *VL, *VR, *h_work, *w1, *w2;
double *w1i, *w2i;
magmaDoubleComplex *w1copy, *w2copy;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double tnrm, result[9];
magma_int_t N, n2, lda, nb, lwork, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double ulp, ulpinv, error;
magma_int_t status = 0;
ulp = lapackf77_dlamch( "P" );
ulpinv = 1./ulp;
magma_opts opts;
parse_opts( argc, argv, &opts );
// need slightly looser bound (60*eps instead of 30*eps) for some tests
opts.tolerance = max( 60., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
// enable at least some minimal checks, if requested
if ( opts.check && !opts.lapack && opts.jobvl == MagmaNoVec && opts.jobvr == MagmaNoVec ) {
fprintf( stderr, "NOTE: Some checks require vectors to be computed;\n"
" set jobvl=V (option -LV), or jobvr=V (option -RV), or both.\n"
" Some checks require running lapack (-l); setting lapack.\n\n");
opts.lapack = true;
}
printf(" N CPU Time (sec) GPU Time (sec) |W_magma - W_lapack| / |W_lapack|\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_dgehrd_nb(N);
lwork = N*(2 + nb);
// generous workspace - required by dget22
lwork = max( lwork, N*(5 + 2*N) );
TESTING_MALLOC_CPU( w1copy, magmaDoubleComplex, N );
TESTING_MALLOC_CPU( w2copy, magmaDoubleComplex, N );
TESTING_MALLOC_CPU( w1, double, N );
TESTING_MALLOC_CPU( w2, double, N );
TESTING_MALLOC_CPU( w1i, double, N );
TESTING_MALLOC_CPU( w2i, double, N );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_PIN( VL, double, n2 );
TESTING_MALLOC_PIN( VR, double, n2 );
TESTING_MALLOC_PIN( h_work, double, lwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dgeev( opts.jobvl, opts.jobvr,
N, h_R, lda, w1, w1i,
VL, lda, VR, lda,
h_work, lwork, opts.queue, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_dgeev returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result
=================================================================== */
if ( opts.check ) {
/* ===================================================================
* Check the result following LAPACK's [zcds]drvev routine.
* The following tests are performed:
* (1) | A * VR - VR * W | / ( n |A| )
*
* Here VR is the matrix of unit right eigenvectors.
* W is a diagonal matrix with diagonal entries W(j).
*
* (2) | |VR(i)| - 1 | and whether largest component real
*
* VR(i) denotes the i-th column of VR.
*
* (3) | A**T * VL - VL * W**T | / ( n |A| )
*
* Here VL is the matrix of unit left eigenvectors, A**T is the
* transpose of A, and W is as above.
*
* (4) | |VL(i)| - 1 | and whether largest component real
*
* VL(i) denotes the i-th column of VL.
*
* (5) W(full) = W(partial, W only) -- currently skipped
* (6) W(full) = W(partial, W and VR)
* (7) W(full) = W(partial, W and VL)
*
* W(full) denotes the eigenvalues computed when both VR and VL
* are also computed, and W(partial) denotes the eigenvalues
* computed when only W, only W and VR, or only W and VL are
* computed.
*
* (8) VR(full) = VR(partial, W and VR)
*
* VR(full) denotes the right eigenvectors computed when both VR
* and VL are computed, and VR(partial) denotes the result
* when only VR is computed.
*
* (9) VL(full) = VL(partial, W and VL)
*
* VL(full) denotes the left eigenvectors computed when both VR
* and VL are also computed, and VL(partial) denotes the result
* when only VL is computed.
*
* (1, 2) only if jobvr = V
* (3, 4) only if jobvl = V
* (5-9) only if check = 2 (option -c2)
================================================================= */
double vmx, vrmx, vtst;
// Initialize result. -1 indicates test was not run.
for( int j = 0; j < 9; ++j )
result[j] = -1.;
if ( opts.jobvr == MagmaVec ) {
// Do test 1: | A * VR - VR * W | / ( n |A| )
// Note this writes result[1] also
lapackf77_dget22( MagmaNoTransStr, MagmaNoTransStr, MagmaNoTransStr,
&N, h_A, &lda, VR, &lda, w1, w1i,
h_work, &result[0] );
result[0] *= ulp;
// Do test 2: | |VR(i)| - 1 | and whether largest component real
result[1] = -1.;
for( int j = 0; j < N; ++j ) {
tnrm = 1.;
if (w1i[j] == 0.)
tnrm = magma_cblas_dnrm2( N, &VR[j*lda], ione );
else if (w1i[j] > 0.)
tnrm = magma_dlapy2( magma_cblas_dnrm2( N, &VR[j*lda], ione ),
magma_cblas_dnrm2( N, &VR[(j+1)*lda], ione ));
result[1] = max( result[1], min( ulpinv, MAGMA_D_ABS(tnrm-1.)/ulp ));
if (w1i[j] > 0.) {
vmx = vrmx = 0.;
for( int jj = 0; jj < N; ++jj ) {
vtst = magma_dlapy2( VR[jj+j*lda], VR[jj+(j+1)*lda]);
if (vtst > vmx)
vmx = vtst;
if ( (VR[jj + (j+1)*lda])==0. &&
MAGMA_D_ABS( VR[jj+j*lda] ) > vrmx)
{
vrmx = MAGMA_D_ABS( VR[jj+j*lda] );
}
}
if (vrmx / vmx < 1. - ulp*2.)
result[1] = ulpinv;
}
}
result[1] *= ulp;
}
if ( opts.jobvl == MagmaVec ) {
// Do test 3: | A**T * VL - VL * W**T | / ( n |A| )
// Note this writes result[3] also
lapackf77_dget22( MagmaTransStr, MagmaNoTransStr, MagmaTransStr,
&N, h_A, &lda, VL, &lda, w1, w1i,
h_work, &result[2] );
result[2] *= ulp;
// Do test 4: | |VL(i)| - 1 | and whether largest component real
result[3] = -1.;
for( int j = 0; j < N; ++j ) {
tnrm = 1.;
if (w1i[j] == 0.)
tnrm = magma_cblas_dnrm2( N, &VL[j*lda], ione );
else if (w1i[j] > 0.)
tnrm = magma_dlapy2( magma_cblas_dnrm2( N, &VL[j*lda], ione ),
magma_cblas_dnrm2( N, &VL[(j+1)*lda], ione ));
result[3] = max( result[3], min( ulpinv, MAGMA_D_ABS(tnrm-1.)/ulp ));
if (w1i[j] > 0.) {
vmx = vrmx = 0.;
for( int jj = 0; jj < N; ++jj ) {
vtst = magma_dlapy2( VL[jj+j*lda], VL[jj+(j+1)*lda]);
if (vtst > vmx)
vmx = vtst;
if ( (VL[jj + (j+1)*lda])==0. &&
MAGMA_D_ABS( VL[jj+j*lda]) > vrmx)
{
vrmx = MAGMA_D_ABS( VL[jj+j*lda] );
}
}
if (vrmx / vmx < 1. - ulp*2.)
result[3] = ulpinv;
}
}
result[3] *= ulp;
}
}
if ( opts.check == 2 ) {
// more extensive tests
// this is really slow because it calls magma_zgeev multiple times
double *LRE, DUM;
TESTING_MALLOC_PIN( LRE, double, n2 );
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
// ----------
// Compute eigenvalues, left and right eigenvectors
magma_dgeev( MagmaVec, MagmaVec,
N, h_R, lda, w1, w1i,
VL, lda, VR, lda,
h_work, lwork, opts.queue, &info );
if (info != 0)
printf("magma_zgeev (case V, V) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// ----------
// Compute eigenvalues only
// These are not exactly equal, and not in the same order, so skip for now.
//lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
//magma_dgeev( MagmaNoVec, MagmaNoVec,
// N, h_R, lda, w2, w2i,
// &DUM, 1, &DUM, 1,
// h_work, lwork, opts.queue, &info );
//if (info != 0)
// printf("magma_dgeev (case N, N) returned error %d: %s.\n",
// (int) info, magma_strerror( info ));
//
//// Do test 5: W(full) = W(partial, W only)
//result[4] = 1;
//for( int j = 0; j < N; ++j )
// if ( w1[j] != w2[j] || w1i[j] != w2i[j] )
// result[4] = 0;
// ----------
// Compute eigenvalues and right eigenvectors
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dgeev( MagmaNoVec, MagmaVec,
N, h_R, lda, w2, w2i,
&DUM, 1, LRE, lda,
h_work, lwork, opts.queue, &info );
if (info != 0)
printf("magma_dgeev (case N, V) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Do test 6: W(full) = W(partial, W and VR)
result[5] = 1;
for( int j = 0; j < N; ++j )
if ( w1[j] != w2[j] || w1i[j] != w2i[j] )
result[5] = 0;
// Do test 8: VR(full) = VR(partial, W and VR)
result[7] = 1;
for( int j = 0; j < N; ++j )
for( int jj = 0; jj < N; ++jj )
if ( ! MAGMA_D_EQUAL( VR[j+jj*lda], LRE[j+jj*lda] ))
result[7] = 0;
// ----------
// Compute eigenvalues and left eigenvectors
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_dgeev( MagmaVec, MagmaNoVec,
N, h_R, lda, w2, w2i,
LRE, lda, &DUM, 1,
h_work, lwork, opts.queue, &info );
if (info != 0)
printf("magma_dgeev (case V, N) returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Do test 7: W(full) = W(partial, W and VL)
result[6] = 1;
for( int j = 0; j < N; ++j )
if ( w1[j] != w2[j] || w1i[j] != w2i[j] )
result[6] = 0;
// Do test 9: VL(full) = VL(partial, W and VL)
result[8] = 1;
for( int j = 0; j < N; ++j )
for( int jj = 0; jj < N; ++jj )
if ( ! MAGMA_D_EQUAL( VL[j+jj*lda], LRE[j+jj*lda] ))
result[8] = 0;
TESTING_FREE_PIN( LRE );
}
/* =====================================================================
Performs operation using LAPACK
Do this after checks, because it overwrites VL and VR.
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgeev( lapack_vec_const(opts.jobvl), lapack_vec_const(opts.jobvr),
&N, h_A, &lda, w2, w2i,
VL, &lda, VR, &lda,
h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_dgeev returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// check | W_magma - W_lapack | / | W |
// need to sort eigenvalues first
// copy them into complex vectors for ease
for( int j=0; j < N; ++j ) {
w1copy[j] = MAGMA_Z_MAKE( w1[j], w1i[j] );
w2copy[j] = MAGMA_Z_MAKE( w2[j], w2i[j] );
}
std::sort( w1copy, &w1copy[N], lessthan );
std::sort( w2copy, &w2copy[N], lessthan );
// adjust sorting to deal with numerical inaccuracy
// search down w2 for eigenvalue that matches w1's eigenvalue
for( int j=0; j < N; ++j ) {
for( int j2=j; j2 < N; ++j2 ) {
magmaDoubleComplex diff = MAGMA_Z_SUB( w1copy[j], w2copy[j2] );
double diff2 = magma_dzlapy2( diff ) / max( magma_dzlapy2( w1copy[j] ), tol );
if ( diff2 < 100*tol ) {
if ( j != j2 ) {
std::swap( w2copy[j], w2copy[j2] );
}
break;
}
}
}
blasf77_zaxpy( &N, &c_neg_one, w2copy, &ione, w1copy, &ione );
error = magma_cblas_dznrm2( N, w1copy, 1 );
error /= magma_cblas_dznrm2( N, w2copy, 1 );
printf("%5d %7.2f %7.2f %8.2e %s\n",
(int) N, cpu_time, gpu_time,
error, (error < tolulp ? "ok" : "failed"));
status += ! (error < tolulp);
}
else {
printf("%5d --- %7.2f\n",
(int) N, gpu_time);
}
if ( opts.check ) {
// -1 indicates test was not run
if ( result[0] != -1 ) {
printf(" | A * VR - VR * W | / ( n |A| ) = %8.2e %s\n", result[0], (result[0] < tol ? "ok" : "failed"));
}
if ( result[1] != -1 ) {
printf(" | |VR(i)| - 1 | = %8.2e %s\n", result[1], (result[1] < tol ? "ok" : "failed"));
}
if ( result[2] != -1 ) {
printf(" | A'* VL - VL * W'| / ( n |A| ) = %8.2e %s\n", result[2], (result[2] < tol ? "ok" : "failed"));
}
if ( result[3] != -1 ) {
printf(" | |VL(i)| - 1 | = %8.2e %s\n", result[3], (result[3] < tol ? "ok" : "failed"));
}
if ( result[4] != -1 ) {
printf(" W (full) == W (partial, W only) %s\n", (result[4] == 1. ? "ok" : "failed"));
}
if ( result[5] != -1 ) {
printf(" W (full) == W (partial, W and VR) %s\n", (result[5] == 1. ? "ok" : "failed"));
}
if ( result[6] != -1 ) {
printf(" W (full) == W (partial, W and VL) %s\n", (result[6] == 1. ? "ok" : "failed"));
}
if ( result[7] != -1 ) {
printf(" VR (full) == VR (partial, W and VR) %s\n", (result[7] == 1. ? "ok" : "failed"));
}
if ( result[8] != -1 ) {
printf(" VL (full) == VL (partial, W and VL) %s\n", (result[8] == 1. ? "ok" : "failed"));
}
int newline = 0;
if ( result[0] != -1 ) {
status += ! (result[0] < tol);
newline = 1;
}
if ( result[1] != -1 ) {
status += ! (result[1] < tol);
newline = 1;
}
if ( result[2] != -1 ) {
status += ! (result[2] < tol);
newline = 1;
}
if ( result[3] != -1 ) {
status += ! (result[3] < tol);
newline = 1;
}
if ( result[4] != -1 ) {
status += ! (result[4] == 1.);
newline = 1;
}
if ( result[5] != -1 ) {
status += ! (result[5] == 1.);
newline = 1;
}
if ( result[6] != -1 ) {
status += ! (result[6] == 1.);
newline = 1;
}
if ( result[7] != -1 ) {
status += ! (result[7] == 1.);
newline = 1;
}
if ( result[8] != -1 ) {
status += ! (result[8] == 1.);
newline = 1;
}
if ( newline ) {
printf( "\n" );
}
}
TESTING_FREE_CPU( w1copy );
TESTING_FREE_CPU( w2copy );
TESTING_FREE_CPU( w1 );
TESTING_FREE_CPU( w2 );
TESTING_FREE_CPU( w1i );
TESTING_FREE_CPU( w2i );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( VL );
TESTING_FREE_PIN( VR );
TESTING_FREE_PIN( h_work );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf
*/
int main( magma_int_t argc, char** argv)
{
cuDoubleComplex *h_A, *h_R, *h_A2, *h_A3, *h_work, *h_work2, *tau, *d_work2;
cuDoubleComplex *d_A, *d_work;
float gpu_perf, cpu_perf, cpu2_perf;
double flops;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2, lda, M=0;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, j, info[1];
magma_int_t loop = argc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t num_cores = 4;
magma_int_t num_gpus = 0;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
else if (strcmp("-C", argv[i])==0)
num_cores = atoi(argv[++i]);
}
if ((M>0 && N>0) || (M==0 && N==0)) {
printf(" testing_zgeqrf_mc -M %d -N %d \n\n", M, N);
if (M==0 && N==0) {
M = N = size[9];
loop = 1;
}
} else {
printf("\nUsage: \n");
printf(" testing_zgeqrf_mc -M %d -N %d -B 128 -T 1\n\n", 1024, 1024);
exit(1);
}
} else {
printf("\nUsage: \n");
printf(" testing_zgeqrf_mc -M %d -N %d -B 128 -T 1\n\n", 1024, 1024);
M = N = size[9];
}
n2 = M * N;
magma_int_t min_mn = min(M,N);
/* Allocate host memory for the matrix */
h_A2 = (cuDoubleComplex*)malloc(n2 * sizeof(h_A2[0]));
if (h_A2 == 0) {
fprintf (stderr, "!!!! host memory allocation error (A2)\n");
}
magma_int_t lwork = n2;
h_work2 = (cuDoubleComplex*)malloc(lwork * sizeof(cuDoubleComplex));
if (h_work2 == 0) {
fprintf (stderr, "!!!! host memory allocation error (h_work2)\n");
}
h_A3 = (cuDoubleComplex*)malloc(n2 * sizeof(h_A3[0]));
if (h_A3 == 0) {
fprintf (stderr, "!!!! host memory allocation error (A3)\n");
}
tau = (cuDoubleComplex*)malloc(min_mn * sizeof(cuDoubleComplex));
if (tau == 0) {
fprintf (stderr, "!!!! host memory allocation error (tau)\n");
}
/* Initialize MAGMA hardware context, seeting how many CPU cores
and how many GPUs to be used in the consequent computations */
magma_context *context;
context = magma_init(NULL, NULL, 0, num_cores, num_gpus, argc, argv);
printf("\n\n");
printf(" M N LAPACK Gflop/s Multi-core Gflop/s ||R||_F / ||A||_F\n");
printf("===========================================================================\n");
for(i=0; i<10; i++){
if (loop == 1) {
M = N = size[i];
n2 = M*N;
}
flops = FLOPS( (double)M, (double)N ) / 1000000;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A2 );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A2, &M, h_A3, &M );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_zgeqrf(&M, &N, h_A3, &M, tau, h_work2, &lwork, info);
end = get_current_time();
if (info[0] < 0)
printf("Argument %d of sgeqrf had an illegal value.\n", -info[0]);
cpu2_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using multicore
=================================================================== */
start = get_current_time();
magma_zgeqrf_mc(context, &M, &N, h_A2, &M, tau, h_work2, &lwork, info);
end = get_current_time();
if (info[0] < 0)
printf("Argument %d of sgeqrf had an illegal value.\n", -info[0]);
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
double work[1], matnorm = 1.;
cuDoubleComplex mone = MAGMA_Z_NEG_ONE;
magma_int_t one = 1;
matnorm = lapackf77_zlange("f", &M, &N, h_A2, &M, work);
blasf77_zaxpy(&n2, &mone, h_A2, &one, h_A3, &one);
printf("%5d %5d %6.2f %6.2f %e\n",
M, N, cpu2_perf, cpu_perf,
lapackf77_zlange("f", &M, &N, h_A3, &M, work) / matnorm);
if (loop != 1)
break;
}
/* Memory clean up */
free(h_A2);
free(tau);
free(h_A3);
free(h_work2);
/* Shut down the MAGMA context */
magma_finalize(context);
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zlaset
Code is very similar to testing_zlacpy.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];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magmaDoubleComplex offdiag = MAGMA_Z_MAKE( 1.2000, 6.7000 );
magmaDoubleComplex diag = MAGMA_Z_MAKE( 3.1415, 2.7183 );
magma_int_t M, N, size, lda, ldb, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("==================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
//M += 2; // space for insets
//N += 2;
lda = M;
ldb = lda;
ldda = ((M+31)/32)*32;
size = lda*N;
if ( uplo[iuplo] == MagmaLower || uplo[iuplo] == MagmaUpper ) {
// save triangle (with diagonal)
// TODO wrong for trapezoid
gbytes = sizeof(magmaDoubleComplex) * 0.5*N*(N+1) / 1e9;
}
else {
// save entire matrix
gbytes = sizeof(magmaDoubleComplex) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, size );
TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, size );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_Z_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
gpu_time = magma_sync_wtime( 0 );
//magmablas_zlaset( uplo[iuplo], M-2, N-2, offdiag, diag, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_zlaset( uplo[iuplo], M, N, offdiag, diag, d_A, ldda );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//magma_int_t M2 = M-2; // inset by 1 row & col
//magma_int_t N2 = N-2;
//lapackf77_zlaset( lapack_uplo_const( uplo[iuplo] ), &M2, &N2, &offdiag, &diag, h_A+1+lda, &lda );
lapackf77_zlaset( lapack_uplo_const( uplo[iuplo] ), &M, &N, &offdiag, &diag, h_A, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_zgetmatrix( M, N, d_A, ldda, h_R, lda );
blasf77_zaxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &M, &N, h_R, &lda, work);
printf("%4c %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
lapacke_uplo_const( uplo[iuplo] ), (int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
ZLATRD2 reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX_16 array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd2(magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dW, magma_int_t lddw,
magmaDoubleComplex *dwork, magma_int_t ldwork)
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
#define dW(i, j) (dW + (j)*lddw + (i))
magma_int_t i;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_zmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside zlatrd
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_zgemv("No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_zgemv("No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_MAKE( 1, 0 );
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1 );
//#if (GPUSHMEM < 200)
//magma_zhemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
// dA(0, i), ione, c_zero, dW(0, iw), ione);
//#else
magmablas_zhemv_work(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( stream );
if (i < n-1) {
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy(&i, &alpha, A(0, i), &ione,
W(0, iw), &ione);
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(i, 0), &ldw);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(i, 0), &ldw);
lapackf77_zlacgv(&i, A(i, 0), &lda);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, A(i, 0), &lda);
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_MAKE( 1, 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
//#if (GPUSHMEM < 200)
//magma_zhemv(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
// dW(i+1, i), ione);
//#else
magmablas_zhemv_work(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
// 3. Here is where we need it
magma_queue_sync( stream );
if (i != 0)
blasf77_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* magma_zlatrd */
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
magmaDoubleComplex *hA, *hR;
magmaDoubleComplex_ptr dA;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] =
{ 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 8160, 8192 };
magma_int_t i, info;
magmaDoubleComplex mz_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], matnorm, diffnorm;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
}
if (N>0) size[0] = size[9] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_zpotrf2_gpu -N %d\n\n", 1024);
}
/* Initialize */
magma_queue_t queue1, queue2;
magma_device_t device;
magma_int_t num = 0;
magma_int_t err;
magma_init();
err = magma_getdevices( &device, 2, &num );
if ( err != 0 or num < 1 ) {
fprintf( stderr, "magma_getdevices failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( device, &queue1 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( device, &queue2 );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
magma_queue_t queues[2] = {queue1, queue2};
/* Allocate memory for the largest matrix */
N = size[9];
n2 = N * N;
ldda = ((N+31)/32) * 32;
TESTING_MALLOC_CPU( hA, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( hR, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*N );
printf("\n\n");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma-R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS( (double)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, hA );
/* Symmetrize and increase the diagonal */
for( int i = 0; i < N; ++i ) {
hA(i,i) = MAGMA_Z_MAKE( MAGMA_Z_REAL(hA(i,i)) + N, 0 );
for( int j = 0; j < i; ++j ) {
hA(i, j) = MAGMA_Z_CNJG( hA(j,i) );
}
}
lapackf77_zlacpy( MagmaFullStr, &N, &N, hA, &lda, hR, &lda );
/* Warm up to measure the performance */
magma_zsetmatrix( N, N, hA, lda, dA, 0, ldda, queue1);
clFinish(queue1);
magma_zpotrf2_gpu( MagmaLower, N, dA, 0, ldda, queues, &info );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( N, N, hA, lda, dA, 0, ldda, queue1 );
clFinish(queue1);
gpu_time = magma_wtime();
magma_zpotrf2_gpu( MagmaLower, N, dA, 0, ldda, queues, &info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf( "magma_zpotrf2 had error %d.\n", info );
gpu_perf = gflops / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( MagmaLowerStr, &N, hA, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf( "lapackf77_zpotrf had error %d.\n", info );
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result compared to LAPACK
|R_magma - R_lapack| / |R_lapack|
=================================================================== */
magma_zgetmatrix( N, N, dA, 0, ldda, hR, lda, queue1 );
matnorm = lapackf77_zlange("f", &N, &N, hA, &lda, work);
blasf77_zaxpy(&n2, &mz_one, hA, &ione, hR, &ione);
diffnorm = lapackf77_zlange("f", &N, &N, hR, &lda, work);
printf( "%5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
N, cpu_perf, cpu_time, gpu_perf, gpu_time, diffnorm / matnorm );
if (argc != 1)
break;
}
/* clean up */
TESTING_FREE_CPU( hA );
TESTING_FREE_PIN( hR );
TESTING_FREE_DEV( dA );
magma_queue_destroy( queue1 );
magma_queue_destroy( queue2 );
magma_finalize();
}
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magma_int_t N, n2, lda, ldda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magmaDoubleComplex **d_A_array = NULL;
magma_int_t *dinfo_magma;
magma_int_t batchCount;
magma_queue_t queue = magma_stream;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
batchCount = opts.batchcount;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("BatchCount N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
ldda = lda = ((N+31)/32)*32;
n2 = lda* N * batchCount;
gflops = batchCount * FLOPS_ZPOTRF( N ) / 1e9 ;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2);
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2);
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda * N * batchCount);
TESTING_MALLOC_DEV( dinfo_magma, magma_int_t, batchCount);
magma_malloc((void**)&d_A_array, batchCount * sizeof(*d_A_array));
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
for(int i=0; i<batchCount; i++)
{
magma_zmake_hpd( N, h_A + i * lda * N, lda );// need modification
}
magma_int_t columns = N * batchCount;
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &(columns), h_A, &lda, h_R, &lda );
magma_zsetmatrix( N, columns, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
zset_pointer(d_A_array, d_A, ldda, 0, 0, ldda * N, batchCount, queue);
gpu_time = magma_sync_wtime(NULL);
info = magma_zpotrf_batched( opts.uplo, N, d_A_array, ldda, dinfo_magma, batchCount, queue);
gpu_time = magma_sync_wtime(NULL) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_int_t *cpu_info = (magma_int_t*) malloc(batchCount*sizeof(magma_int_t));
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_zpotrf_batched matrix %d returned internal error %d\n",i, (int)cpu_info[i] );
}
}
if (info != 0)
printf("magma_zpotrf_batched returned argument error %d: %s.\n", (int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for(int i=0; i<batchCount; i++)
{
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A + i * lda * N, &lda, &info );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, columns, d_A, ldda, h_R, lda );
magma_int_t NN = lda*N;
char const uplo = 'l'; // lapack_uplo_const(opts.uplo)
double err = 0.0;
for(int i=0; i<batchCount; i++)
{
error = lapackf77_zlanhe("f", &uplo, &N, h_A + i * lda*N, &lda, work);
blasf77_zaxpy(&NN, &c_neg_one, h_A + i * lda*N, &ione, h_R + i * lda*N, &ione);
error = lapackf77_zlanhe("f", &uplo, &N, h_R + i * lda*N, &lda, work) / error;
if ( isnan(error) || isinf(error) ) {
err = error;
break;
}
err = max(fabs(error),err);
}
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int)batchCount, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., err, (error < tol ? "ok" : "failed"));
status += ! (err < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int)batchCount, (int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_A_array );
TESTING_FREE_DEV( dinfo_magma );
free(cpu_info);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, ldda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("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;
ldda = ((N+31)/32)*32;
gflops = FLOPS_ZPOTRI( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* factorize matrix */
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
magma_zpotrf_gpu( opts.uplo, N, d_A, ldda, &info );
// check for exact singularity
//magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
//h_R[ 10 + 10*lda ] = MAGMA_Z_MAKE( 0.0, 0.0 );
//magma_zsetmatrix( N, N, h_R, lda, d_A, ldda );
gpu_time = magma_wtime();
magma_zpotri_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_zpotri( 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_zpotri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
error = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %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 ztrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time=0, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double magma_error, cublas_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *ipiv;
magmaDoubleComplex *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2;
magmaDoubleComplex *d_A, *d_B;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("side = %s, uplo = %s, transA = %s, diag = %s \n",
lapack_side_const(opts.side), lapack_uplo_const(opts.uplo),
lapack_trans_const(opts.transA), lapack_diag_const(opts.diag) );
printf(" M N 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];
gflops = FLOPS_ZTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC_CPU( h_B, magmaDoubleComplex, ldb*N );
TESTING_MALLOC_CPU( h_B1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC_CPU( h_X1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC_CPU( h_X2, magmaDoubleComplex, ldb*N );
TESTING_MALLOC_CPU( h_Bcublas, magmaDoubleComplex, ldb*N );
TESTING_MALLOC_CPU( h_Bmagma, magmaDoubleComplex, ldb*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, Ak );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_B, magmaDoubleComplex, lddb*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( &Ak, &Ak, h_A, &lda, ipiv, &info );
for( int j = 0; j < Ak; ++j ) {
for( int i = 0; i < j; ++i ) {
*h_A(i,j) = *h_A(j,i);
}
}
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(magmaDoubleComplex));
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_zsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_ztrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasZtrsm( handle, cublas_side_const(opts.side), cublas_uplo_const(opts.uplo),
cublas_trans_const(opts.transA), cublas_diag_const(opts.diag),
M, N,
&alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ztrsm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(magmaDoubleComplex));
magmaDoubleComplex alpha2 = MAGMA_Z_DIV( c_one, alpha );
blasf77_ztrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
double norm1 = lapackf77_zlange( "M", &M, &N, h_X1, &ldb, work );
double normx = lapackf77_zlange( "M", &M, &N, h_Bmagma, &ldb, work );
double normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(magmaDoubleComplex));
blasf77_ztrmm( lapack_side_const(opts.side), lapack_uplo_const(opts.uplo), lapack_trans_const(opts.transA), lapack_diag_const(opts.diag),
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_zlange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_zlange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_B1 );
TESTING_FREE_CPU( h_X1 );
TESTING_FREE_CPU( h_X2 );
TESTING_FREE_CPU( h_Bcublas );
TESTING_FREE_CPU( h_Bmagma );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotf2_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex_ptr d_A;
magma_int_t N, n2, lda, ldda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double Anorm, error, work[1];
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% N CPU Gflop/s (ms) GPU Gflop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("%%=======================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
gflops = FLOPS_ZPOTRF( N ) / 1e9;
if ( N > 512 ) {
printf( "%5d skipping because zpotf2 does not support N > 512\n", (int) N );
continue;
}
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( opts.queue );
magma_zpotf2_gpu( opts.uplo, N, d_A, ldda, opts.queue, &info );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_zpotf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( 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_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda, opts.queue );
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
Anorm = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
error = lapackf77_zlange("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*1000., gpu_perf, gpu_time*1000.,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zlarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t M, N, K, size, ldc, ldv, ldt, ldw, nv;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double error, work[1];
magma_int_t status = 0;
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans [] = { MagmaConjTrans, MagmaNoTrans };
magma_direct_t direct[] = { MagmaForward, MagmaBackward };
magma_storev_t storev[] = { MagmaColumnwise, MagmaRowwise };
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
if ( M < K || N < K || K <= 0 ) {
printf( "%5d %5d %5d skipping because zlarfb requires M >= K, N >= K, K >= 0\n",
(int) M, (int) N, (int) K );
continue;
}
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
ldc = ((M+31)/32)*32;
ldt = ((K+31)/32)*32;
ldw = (side[iside] == MagmaLeft ? N : M);
// (ldv, nv) get swapped later if rowwise
ldv = (side[iside] == MagmaLeft ? M : N);
nv = K;
// Allocate memory for matrices
magmaDoubleComplex *C, *R, *V, *T, *W;
TESTING_MALLOC_CPU( C, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_CPU( R, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_CPU( V, magmaDoubleComplex, ldv*K );
TESTING_MALLOC_CPU( T, magmaDoubleComplex, ldt*K );
TESTING_MALLOC_CPU( W, magmaDoubleComplex, ldw*K );
magmaDoubleComplex *dC, *dV, *dT, *dW;
TESTING_MALLOC_DEV( dC, magmaDoubleComplex, ldc*N );
TESTING_MALLOC_DEV( dV, magmaDoubleComplex, ldv*K );
TESTING_MALLOC_DEV( dT, magmaDoubleComplex, ldt*K );
TESTING_MALLOC_DEV( dW, magmaDoubleComplex, ldw*K );
// C is M x N.
size = ldc*N;
lapackf77_zlarnv( &ione, ISEED, &size, C );
//printf( "C=" ); magma_zprint( M, N, C, ldc );
// V is ldv x nv. See larfb docs for description.
// if column-wise and left, M x K
// if column-wise and right, N x K
// if row-wise and left, K x M
// if row-wise and right, K x N
size = ldv*nv;
lapackf77_zlarnv( &ione, ISEED, &size, V );
if ( storev[istor] == MagmaColumnwise ) {
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_zlaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, &V[(ldv-K)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_zlaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, &V[(nv-K)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_zprint( ldv, nv, V, ldv );
// T is K x K, upper triangular for forward, and lower triangular for backward
magma_int_t k1 = K-1;
size = ldt*K;
lapackf77_zlarnv( &ione, ISEED, &size, T );
if ( direct[idir] == MagmaForward ) {
lapackf77_zlaset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_zlaset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_zprint( K, K, T, ldt );
magma_zsetmatrix( M, N, C, ldc, dC, ldc );
magma_zsetmatrix( ldv, nv, V, ldv, dV, ldv );
magma_zsetmatrix( K, K, T, ldt, dT, ldt );
lapackf77_zlarfb( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
lapack_direct_const( direct[idir] ), lapack_storev_const( storev[istor] ),
&M, &N, &K,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_zprint( M, N, C, ldc );
magma_zlarfb_gpu( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw );
magma_zgetmatrix( M, N, dC, ldc, R, ldc );
//printf( "dHC=" ); magma_zprint( M, N, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
error = lapackf77_zlange( "Fro", &M, &N, C, &ldc, work );
size = ldc*N;
blasf77_zaxpy( &size, &c_neg_one, C, &ione, R, &ione );
error = lapackf77_zlange( "Fro", &M, &N, R, &ldc, work ) / error;
printf( "%5d %5d %5d %c %c %c %c %8.2e %s\n",
(int) M, (int) N, (int) K,
lapacke_storev_const(storev[istor]), lapacke_side_const(side[iside]),
lapacke_direct_const(direct[idir]), lapacke_trans_const(trans[itran]),
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( V );
TESTING_FREE_CPU( T );
TESTING_FREE_CPU( W );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dV );
TESTING_FREE_DEV( dT );
TESTING_FREE_DEV( dW );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}}}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zungqr_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *hA, *hR, *tau, *h_work;
magmaDoubleComplex *dA, *dT;
magma_int_t m, n, k;
magma_int_t n2, lda, ldda, lwork, min_mn, nb, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" m n k CPU GFlop/s (sec) GPU GFlop/s (sec) ||R|| / ||A||\n");
printf("=========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
if ( m < n || n < k ) {
printf( "skipping m %d, n %d, k %d because m < n or n < k\n", (int) m, (int) n, (int) k );
continue;
}
lda = m;
ldda = ((m + 31)/32)*32;
n2 = lda*n;
min_mn = min(m, n);
nb = magma_get_zgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_ZUNGQR( m, n, k ) / 1e9;
TESTING_MALLOC_PIN( hA, magmaDoubleComplex, lda*n );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_CPU( hR, magmaDoubleComplex, lda*n );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*n );
TESTING_MALLOC_DEV( dT, magmaDoubleComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_zlarnv( &ione, ISEED, &n2, hA );
lapackf77_zlacpy( MagmaUpperLowerStr, &m, &n, hA, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( m, n, hA, lda, dA, ldda );
magma_zgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_zgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
gpu_time = magma_wtime();
magma_zungqr_gpu( m, n, k, dA, ldda, tau, dT, nb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get dA back to the CPU to compare with the CPU result.
magma_zgetmatrix( m, n, dA, ldda, hR, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
error = lapackf77_zlange("f", &m, &n, hA, &lda, work );
lapackf77_zgeqrf( &m, &n, hA, &lda, tau, h_work, &lwork, &info );
if (info != 0)
printf("lapackf77_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
cpu_time = magma_wtime();
lapackf77_zungqr( &m, &n, &k, hA, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_zaxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_zlange("f", &m, &n, hR, &lda, work) / error;
printf("%5d %5d %5d %7.1f (%7.2f) %7.1f (%7.2f) %8.2e\n",
(int) m, (int) n, (int) k,
cpu_perf, cpu_time, gpu_perf, gpu_time, error );
}
else {
printf("%5d %5d %5d --- ( --- ) %7.1f (%7.2f) --- \n",
(int) m, (int) n, (int) k,
gpu_perf, gpu_time );
}
TESTING_FREE_PIN( hA );
TESTING_FREE_PIN( h_work );
TESTING_FREE_CPU( hR );
TESTING_FREE_CPU( tau );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dT );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_x, *h_x2, *h_tau, *h_tau2;
magmaDoubleComplex *d_x, *d_tau;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double error, error2, work[1];
magma_int_t N, nb, lda, ldda, size;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}; magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
// does larfg on nb columns, one after another
nb = (opts.nb > 0 ? opts.nb : 64);
magma_queue_t queue = 0;
printf(" N nb CPU GFLop/s (ms) GPU GFlop/s (ms) error tau error\n");
printf("==========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_ZLARFG( N ) / 1e9 * nb;
TESTING_MALLOC_CPU( h_x, magmaDoubleComplex, N*nb );
TESTING_MALLOC_CPU( h_x2, magmaDoubleComplex, N*nb );
TESTING_MALLOC_CPU( h_tau, magmaDoubleComplex, nb );
TESTING_MALLOC_CPU( h_tau2, magmaDoubleComplex, nb );
TESTING_MALLOC_DEV( d_x, magmaDoubleComplex, ldda*nb );
TESTING_MALLOC_DEV( d_tau, magmaDoubleComplex, nb );
/* Initialize the vectors */
size = N*nb;
lapackf77_zlarnv( &ione, ISEED, &size, h_x );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_zsetmatrix( N, nb, h_x, N, d_x, ldda );
gpu_time = magma_sync_wtime( queue );
for( int j = 0; j < nb; ++j ) {
magmablas_zlarfg( N, &d_x[0+j*ldda], &d_x[1+j*ldda], ione, &d_tau[j] );
}
gpu_time = magma_sync_wtime( queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_zgetmatrix( N, nb, d_x, ldda, h_x2, N );
magma_zgetvector( nb, d_tau, 1, h_tau2, 1 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
for( int j = 0; j < nb; ++j ) {
lapackf77_zlarfg( &N, &h_x[0+j*lda], &h_x[1+j*lda], &ione, &h_tau[j] );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Error Computation and Performance Comparison
=================================================================== */
blasf77_zaxpy( &size, &c_neg_one, h_x, &ione, h_x2, &ione );
error = lapackf77_zlange( "F", &N, &nb, h_x2, &N, work )
/ lapackf77_zlange( "F", &N, &nb, h_x, &N, work );
// tau can be 0
blasf77_zaxpy( &nb, &c_neg_one, h_tau, &ione, h_tau2, &ione );
error2 = lapackf77_zlange( "F", &nb, &ione, h_tau, &nb, work );
if ( error2 != 0 ) {
error2 = lapackf77_zlange( "F", &nb, &ione, h_tau2, &nb, work ) / error2;
}
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) N, (int) nb, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error, error2,
(error < tol && error2 < tol ? "ok" : "failed") );
status += ! (error < tol && error2 < tol);
TESTING_FREE_CPU( h_x );
TESTING_FREE_CPU( h_x2 );
TESTING_FREE_CPU( h_tau );
TESTING_FREE_DEV( d_x );
TESTING_FREE_DEV( d_tau );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgelqf_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_A;
magma_int_t M, N, n2, lda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
nb = magma_get_zgeqrf_nb(M);
gflops = FLOPS_ZGELQF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgelqf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, M*nb );
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, lda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, N, h_R, lda, d_A, lda );
gpu_time = magma_wtime();
magma_zgelqf_gpu( M, N, d_A, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgelqf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgelqf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_zgelqf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( M, N, d_A, lda, h_R, lda );
error = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
cuDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
/* Matrix size */
magma_int_t M = 0, N = 0, n2, lda, lwork;
const int MAXTESTS = 10;
magma_int_t msize[MAXTESTS] = { 1024, 2048, 3072, 4032, 5184, 6016, 7040, 8064, 9088, 10112 };
magma_int_t nsize[MAXTESTS] = { 1024, 2048, 3072, 4032, 5184, 6016, 7040, 8064, 9088, 10112 };
magma_int_t i, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t checkres;
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
// process command line arguments
printf( "\nUsage: %s -N <m,n> -c\n", argv[0] );
printf( " -N can be repeated up to %d times. If only m is given, then m=n.\n", MAXTESTS );
printf( " -c or setting $MAGMA_TESTINGS_CHECK runs LAPACK and checks result.\n\n" );
int ntest = 0;
for( int i = 1; i < argc; ++i ) {
if ( strcmp("-N", argv[i]) == 0 && i+1 < argc ) {
magma_assert( ntest < MAXTESTS, "error: -N repeated more than maximum %d tests\n", MAXTESTS );
int m, n;
info = sscanf( argv[++i], "%d,%d", &m, &n );
if ( info == 2 && m > 0 && n > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = n;
}
else if ( info == 1 && m > 0 ) {
msize[ ntest ] = m;
nsize[ ntest ] = m; // implicitly
}
else {
printf( "error: -N %s is invalid; ensure m > 0, n > 0.\n", argv[i] );
exit(1);
}
M = max( M, msize[ ntest ] );
N = max( N, nsize[ ntest ] );
ntest++;
}
else if ( strcmp("-M", argv[i]) == 0 ) {
printf( "-M has been replaced in favor of -N m,n to allow -N to be repeated.\n\n" );
exit(1);
}
else if ( strcmp("-c", argv[i]) == 0 ) {
checkres = true;
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( ntest == 0 ) {
ntest = MAXTESTS;
M = msize[ntest-1];
N = nsize[ntest-1];
}
n2 = M * N;
min_mn = min(M, N);
nb = magma_get_zgeqrf_nb(M);
/* Allocate memory for the matrix */
TESTING_MALLOC( tau, cuDoubleComplex, min_mn );
TESTING_MALLOC( h_A, cuDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, cuDoubleComplex, n2 );
lwork = -1;
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, N*nb );
TESTING_MALLOC( h_work, cuDoubleComplex, lwork );
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( i = 0; i < ntest; ++i ) {
M = msize[i];
N = nsize[i];
min_mn= min(M, N);
lda = M;
n2 = lda*N;
gflops = FLOPS_ZGEQRF( M, N ) / 1e9;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgeqrf(M, N, h_R, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgeqrf returned error %d.\n", (int) info);
if ( checkres ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgeqrf returned error %d.\n", (int) info);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
/* Memory clean up */
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_CUDA_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotrf_mgpu
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_setdevice(0);
magma_timestr_t start, end;
double flops, gpu_perf, cpu_perf;
cuDoubleComplex *h_A, *h_R;
cuDoubleComplex *d_lA[4];
magma_int_t N = 0, n2, mb, nb, nk, lda, ldda, n_local, ldn_local;
//magma_int_t size[10] = {1000,2000,3000,4000,5000,6000,7000,8000,9000,10000};
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t n_sizes = 10, flag = 0;
magma_int_t i, j, k, info, num_gpus0 = 1, num_gpus;
const char *uplo = MagmaLowerStr;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], matnorm;
N = size[n_sizes-1];
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0) {
flag = 1;
N = atoi(argv[++i]);
size[0] = size[n_sizes-1] = N;
}
if (strcmp("-NGPU", argv[i])==0)
num_gpus0 = atoi(argv[++i]);
if (strcmp("-UPLO",argv[i])==0) {
if (strcmp("L",argv[++i])==0) uplo = MagmaLowerStr;
else uplo = MagmaUpperStr;
}
}
if (strcmp(uplo,MagmaLowerStr)==0)
printf("\n testing_zpotrf_mgpu -N %d -NGPU %d -UPLO L\n\n", (int) N, (int) num_gpus0 );
else
printf("\n testing_zpotrf_mgpu -N %d -NGPU %d -UPLO U\n\n", (int) N, (int) num_gpus0 );
} else {
printf("\nDefault: \n");
printf(" testing_zpotrf_mgpu -N %d:%d -NGPU %d -UPLO L\n\n", (int) size[0], (int) size[n_sizes-1], (int) num_gpus0 );
}
if( N <= 0 || num_gpus0 <= 0 ) {
printf( " invalid input N=%d NGPU=%d\n", (int) N, (int) num_gpus0 );
exit(1);
}
/* looking for max. ldda */
ldda = 0;
n2 = 0;
for(i=0; i<n_sizes; i++){
N = size[i];
nb = magma_get_zpotrf_nb(N);
mb = nb;
if( num_gpus0 > N/nb ) {
num_gpus = N/nb;
if( N%nb != 0 ) num_gpus ++;
} else {
num_gpus = num_gpus0;
}
n_local = nb*(1+N/(nb*num_gpus)) * mb*((N+mb-1)/mb);
if( n_local > ldda ) ldda = n_local;
if( n2 < N*N ) n2 = N*N;
if (flag != 0) break;
}
/* Allocate host memory for the matrix */
TESTING_HOSTALLOC( h_A, cuDoubleComplex, n2);
TESTING_HOSTALLOC( h_R, cuDoubleComplex, n2);
/* allocate local matrix on GPU */
for(i=0; i<num_gpus0; i++){
magma_setdevice(i);
TESTING_DEVALLOC( d_lA[i], cuDoubleComplex, ldda );
}
magma_setdevice(0);
printf(" N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("========================================================\n");
for(i=0; i<n_sizes; i++){
N = size[i];
lda = N;
n2 = lda*N;
flops = FLOPS( (double)N ) / 1000000;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_Z_SET2REAL( h_A[i*lda+i], ( MAGMA_Z_REAL(h_A[i*lda+i]) + 1.*N ) );
for(j=0; j<i; j++)
h_A[i*lda+j] = cuConj(h_A[j*lda+i]);
}
}
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
nb = magma_get_zpotrf_nb(N);
if( num_gpus0 > N/nb ) {
num_gpus = N/nb;
if( N%nb != 0 ) num_gpus ++;
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) num_gpus );
} else {
num_gpus = num_gpus0;
}
/* distribute matrix to gpus */
if( lapackf77_lsame(uplo, "U") ) {
/* going through each block-column */
ldda = ((N+mb-1)/mb)*mb;
for(j=0; j<N; j+=nb){
k = (j/nb)%num_gpus;
magma_setdevice(k);
nk = min(nb, N-j);
magma_zsetmatrix( N, nk,
h_A+j*lda, lda,
d_lA[k]+j/(nb*num_gpus)*nb*ldda, ldda );
}
} else {
/* going through each block-row */
ldda = (1+N/(nb*num_gpus))*nb;
for(j=0; j<N; j+=nb){
k = (j/nb)%num_gpus;
magma_setdevice(k);
nk = min(nb, N-j);
magma_zsetmatrix( nk, N,
h_A+j, lda,
d_lA[k]+j/(nb*num_gpus)*nb, ldda );
}
}
magma_setdevice(0);
/* call magma_zpotrf_mgpu */
start = get_current_time();
magma_zpotrf_mgpu(num_gpus, uplo[0], N, d_lA, ldda, &info);
end = get_current_time();
if (info < 0) {
printf("Argument %d of magma_zpotrf_mgpu had an illegal value.\n", (int) -info);
break;
} else if (info != 0) {
printf("magma_zpotrf_mgpu returned info=%d\n", (int) info );
break;
}
gpu_perf = flops / GetTimerValue(start, end);
/* gather matrix from gpus */
if( lapackf77_lsame(uplo, "U") ) {
for(j=0; j<N; j+=nb){
k = (j/nb)%num_gpus;
magma_setdevice(k);
nk = min(nb, N-j);
magma_zgetmatrix( N, nk,
d_lA[k]+j/(nb*num_gpus)*nb*ldda, ldda,
h_R+j*lda, lda );
}
} else {
for(j=0; j<N; j+=nb){
k = (j/nb)%num_gpus;
magma_setdevice(k);
nk = min(nb, N-j);
magma_zgetmatrix( nk, N,
d_lA[k]+j/(nb*num_gpus)*nb, ldda,
h_R+j, lda );
}
}
magma_setdevice(0);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_zpotrf(uplo, &N, h_A, &lda, &info);
end = get_current_time();
if (info < 0) {
printf("Argument %d of zpotrf had an illegal value.\n", (int) -info);
break;
} else if (info != 0) {
printf("lapackf77_zpotrf returned info=%d\n", (int) info );
break;
}
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
matnorm = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %6.2f %6.2f %e\n",
(int) size[i], cpu_perf, gpu_perf,
lapackf77_zlange("f", &N, &N, h_R, &lda, work) / matnorm);
if (flag != 0) break;
}
/* Memory clean up */
TESTING_HOSTFREE( h_A );
TESTING_HOSTFREE( h_R );
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
TESTING_DEVFREE( d_lA[i] );
}
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrmv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cublas_perf, cublas_time, cpu_perf, cpu_time;
double cublas_error, Cnorm, work[1];
magma_int_t N;
magma_int_t Ak;
magma_int_t sizeA;
magma_int_t lda, ldda;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex *h_A, *h_x, *h_xcublas;
magmaDoubleComplex_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 );
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, 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_ZTRMM(opts.side, N, 1) / 1e9;
lda = N;
Ak = N;
ldda = ((lda+31)/32)*32;
sizeA = lda*Ak;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC_CPU( h_x, magmaDoubleComplex, N );
TESTING_MALLOC_CPU( h_xcublas, magmaDoubleComplex, N );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_MALLOC_DEV( d_x, magmaDoubleComplex, N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &N, h_x );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_zsetvector( N, h_x, 1, d_x, 1 );
cublas_time = magma_sync_wtime( NULL );
cublasZtrmv( opts.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_ztrmv( 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
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_zlange( "M", &N, &ione, h_x, &N, work );
blasf77_zaxpy( &N, &c_neg_one, h_x, &ione, h_xcublas, &ione );
cublas_error = lapackf77_zlange( "M", &N, &ione, h_xcublas, &N, work ) / Cnorm;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %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) --- ( --- ) --- ---\n",
(int) N,
cublas_perf, 1000.*cublas_time);
}
TESTING_FREE_CPU( h_A );
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;
}
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;
}
int main(int argc, char **argv)
{
TESTING_INIT();
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magma_int_t ione = 1;
real_Double_t atomics_perf=0, atomics_time=0;
real_Double_t gflops, magma_perf=0, magma_time=0, cublas_perf, cublas_time, cpu_perf, cpu_time;
double magma_error=0, atomics_error=0, cublas_error, work[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 alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Yatomics, *Ycublas, *Ymagma;
magmaDoubleComplex_ptr dA, dX, dY, dwork;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% N MAGMA Gflop/s (ms) Atomics Gflop/s CUBLAS Gflop/s CPU Gflop/s MAGMA error Atomics CUBLAS\n");
printf("%%=====================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_ZHEMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Yatomics, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ycublas, 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 = magma_ceildiv( N, 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 );
// should not use data from the opposite triangle -- fill with NAN to check
magma_int_t N1 = N-1;
if ( opts.uplo == MagmaUpper ) {
lapackf77_zlaset( "Lower", &N1, &N1, &MAGMA_Z_NAN, &MAGMA_Z_NAN, &A[1], &lda );
}
else {
lapackf77_zlaset( "Upper", &N1, &N1, &MAGMA_Z_NAN, &MAGMA_Z_NAN, &A[lda], &lda );
}
lapackf77_zlarnv( &ione, ISEED, &sizeX, X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( N, N, A, lda, dA, ldda );
magma_zsetvector( N, X, incx, dX, incx );
magma_zsetvector( N, Y, incy, dY, incy );
magmablasSetKernelStream( opts.queue ); // opts.handle also uses opts.queue
cublas_time = magma_sync_wtime( opts.queue );
#ifdef HAVE_CUBLAS
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
#else
magma_zhemv( opts.uplo, N, alpha, dA, 0, ldda, dX, 0, incx, beta, dY, 0, incy, opts.queue );
#endif
cublas_time = magma_sync_wtime( opts.queue ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetvector( N, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using CUBLAS - using atomics
=================================================================== */
#ifdef HAVE_CUBLAS
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_ALLOWED );
magma_zsetvector( N, Y, incy, dY, incy );
// sync on queue doesn't work -- need device sync or use NULL stream -- bug in CUBLAS?
atomics_time = magma_sync_wtime( NULL /*opts.queue*/ );
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
atomics_time = magma_sync_wtime( NULL /*opts.queue*/ ) - atomics_time;
atomics_perf = gflops / atomics_time;
magma_zgetvector( N, dY, incy, Yatomics, incy );
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_NOT_ALLOWED );
#endif
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
#ifdef HAVE_CUBLAS
magma_zsetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( opts.queue );
if ( opts.version == 1 ) {
magmablas_zhemv_work( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy, dwork, ldwork, opts.queue );
}
else {
// non-work interface (has added overhead)
magmablas_zhemv( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy );
}
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( N, dY, incy, Ymagma, incy );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_zhemv( 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, Ycublas, &incy );
cublas_error = lapackf77_zlange( "M", &N, &ione, Ycublas, &N, work ) / N;
#ifdef HAVE_CUBLAS
blasf77_zaxpy( &N, &c_neg_one, Y, &incy, Yatomics, &incy );
atomics_error = lapackf77_zlange( "M", &N, &ione, Yatomics, &N, work ) / N;
blasf77_zaxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_zlange( "M", &N, &ione, Ymagma, &N, work ) / N;
#endif
bool okay = (magma_error < tol && cublas_error < tol && atomics_error < tol);
status += ! okay;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e %s\n",
(int) N,
magma_perf, 1000.*magma_time,
atomics_perf, 1000.*atomics_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error, atomics_error,
(okay ? "ok" : "failed"));
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Yatomics );
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" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrtri
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t N, n2, lda, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double 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)
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||_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_ZTRTRI( N ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
/* ====================================================================
Initialize the matrix
=================================================================== */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
if ( opts.warmup ) {
magma_zpotrf( opts.uplo, N, h_R, lda, &info );
magma_ztrtri( opts.uplo, opts.diag, N, h_R, lda, &info );
lapackf77_zlacpy( MagmaFullStr, &N, &N, h_A, &lda, h_R, &lda );
}
/* factorize matrix */
magma_zpotrf( opts.uplo, N, h_R, lda, &info );
// check for exact singularity
//h_R[ 10 + 10*lda ] = MAGMA_Z_ZERO;
gpu_time = magma_wtime();
magma_ztrtri( opts.uplo, opts.diag, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_ztrtri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
lapackf77_zpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime();
lapackf77_ztrtri( lapack_uplo_const(opts.uplo), lapack_diag_const(opts.diag), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_ztrtri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
Anorm = lapackf77_zlantr("f", lapack_uplo_const(opts.uplo), MagmaNonUnitStr, &N, &N, h_A, &lda, work);
error = lapackf77_zlantr("f", lapack_uplo_const(opts.uplo), MagmaNonUnitStr, &N, &N, h_R, &lda, work) / Anorm;
bool okay = (error < tol);
status += ! okay;
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, (okay ? "ok" : "failed") );
}
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 ztranspose
Code is very similar to testing_zsymmetrize.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;
double error, error2, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_B, *h_R;
magmaDoubleComplex *d_A, *d_B;
magma_int_t M, N, size, lda, ldda, ldb, lddb;
magma_int_t ione = 1;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("Inplace transpose requires M==N.\n");
printf(" M N CPU GByte/s (sec) GPU GByte/s (sec) check Inplace GB/s (sec) check\n");
printf("====================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
lda = M;
ldda = ((M+31)/32)*32;
ldb = N;
lddb = ((N+31)/32)*32;
// load entire matrix, save entire matrix
gbytes = sizeof(magmaDoubleComplex) * 2.*M*N / 1e9;
// input is M x N
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
// output is N x M
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*M );
TESTING_MALLOC( h_R, magmaDoubleComplex, ldb*M );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*M );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_Z_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_Z_MAKE( i + j/10000., j );
}
}
magma_zsetmatrix( 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_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( N, M, h_B, ldb, d_B, lddb );
gpu_time = magma_sync_wtime( 0 );
//magmablas_ztranspose2( d_B+1+lddb, lddb, d_A+1+ldda, ldda, M-2, N-2 ); // inset by 1 row & col
magmablas_ztranspose2( d_B, lddb, d_A, ldda, M, N );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* ====================================================================
Performs operation using MAGMA, in-place
=================================================================== */
if ( M == N ) {
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 );
//magmablas_ztranspose_inplace( N-2, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_ztranspose_inplace( N, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 ) - gpu_time2;
gpu_perf2 = gbytes / gpu_time2;
}
/* =====================================================================
Check the result
=================================================================== */
size = ldb*M;
magma_zgetmatrix( N, M, d_B, lddb, h_R, ldb );
blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_zlange("f", &N, &M, h_R, &ldb, work );
if ( M == N ) {
magma_zgetmatrix( N, M, d_A, ldda, h_R, ldb );
blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error2 = lapackf77_zlange("f", &N, &M, h_R, &ldb, work );
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %4s %7.2f (%7.2f) %4s\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
(error == 0. ? "ok" : "failed"),
gpu_perf2, gpu_time2,
(error2 == 0. ? "ok" : "failed") );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %4s --- ( --- )\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
(error == 0. ? "ok" : "failed") );
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( h_R );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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 zgeqrf_mgpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_lA[ MagmaMaxGPUs ];
magma_int_t M, N, n2, lda, ldda, n_local, ngpu;
magma_int_t info, min_mn, nb, lhwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= (opts.check == 2); // check (-c2) implies lapack (-l)
magma_int_t status = 0;
double tol, eps = lapackf77_dlamch("E");
tol = opts.tolerance * eps;
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("================================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F /(M*||A||_F)\n");
printf("==========================================================================\n");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_zgeqrf_nb( M );
gflops = FLOPS_ZGEQRF( M, N ) / 1e9;
// ngpu must be at least the number of blocks
ngpu = min( opts.ngpu, int((N+nb-1)/nb) );
if ( ngpu < opts.ngpu ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
}
// query for workspace size
lhwork = -1;
lapackf77_zgeqrf( &M, &N, NULL, &M, NULL, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_Z_REAL( tmp[0] );
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_work, magmaDoubleComplex, lhwork );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
// Allocate device memory
for( int dev = 0; dev < ngpu; dev++ ) {
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaDoubleComplex, ldda*n_local );
}
/* Initialize the matrix */
for ( int j=0; j<4; j++ )
ISEED2[j] = ISEED[j]; // save seeds
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
magmaDoubleComplex *tau2;
TESTING_MALLOC_CPU( tau2, magmaDoubleComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_zgeqrf( &M, &N, h_A, &M, tau2, h_work, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau2 );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( M, N, h_R, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_zgeqrf2_mgpu( ngpu, M, N, d_lA, ldda, tau, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgeqrf2 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_zgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_R, lda, ngpu, nb );
magma_queue_sync( NULL );
if ( opts.check == 1 && M >= N ) {
/* =====================================================================
Check the result -- zqrt02 requires M >= N
=================================================================== */
magma_int_t lwork = n2+N;
magmaDoubleComplex *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC_CPU( h_W1, magmaDoubleComplex, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, magmaDoubleComplex, n2 ); // R
TESTING_MALLOC_CPU( h_W3, magmaDoubleComplex, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, double, M ); // RWORK
lapackf77_zlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_zqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0], results[1] );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0], results[1] );
}
// todo also check results[1] < tol?
printf(" %s\n", (results[0] < tol ? "ok" : "failed"));
status += ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
}
else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_zlange("f", &M, &N, h_A, &lda, work );
blasf77_zaxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_zlange("f", &M, &N, h_R, &lda, work ) / (min_mn*error);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) ---",
(int) M, (int) N, gpu_perf, gpu_time);
}
printf("%s\n", (opts.check != 0 ? " (error check only for M >= N)" : ""));
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
for( int dev=0; dev < ngpu; dev++ ){
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
