/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvdx
*/
int main( int argc, char** argv)
{
//#define USE_MGPU
#ifdef USE_MGPU
TESTING_CUDA_INIT_MGPU();
#else
TESTING_CUDA_INIT();
#endif
magma_int_t nrgpu =1;
cuDoubleComplex *h_A, *h_R, *h_B, *h_S, *h_work;
double *rwork, *w1, *w2;
magma_int_t *iwork;
double gpu_time, cpu_time;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t size[4] = {1024,2048,4100,6001};
magma_int_t i, itype, info;
magma_int_t ione = 1, izero = 0;
magma_int_t five = 5;
cuDoubleComplex c_zero = MAGMA_Z_ZERO;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double d_one = 1.;
double d_neg_one = -1.;
double d_ten = 10.;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t il,iu,m1,m2;
double vl,vu;
double fraction_ev = 0;
//const char *uplo = MagmaLowerStr;
char *uplo = (char*)MagmaLowerStr;
//char *uplo = (char*)MagmaUpperStr;
char *jobz = (char*)MagmaVectorsStr;
char range = 'A';
itype = 1;
magma_int_t checkres;
double result[2];
int flagN = 0;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0){
N = atoi(argv[++i]);
if (N>0){
printf(" testing_zhegvdx -N %d\n\n", (int) N);
flagN=1;
}
else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -N %d\n\n", (int) N);
exit(1);
}
}
if (strcmp("-ngpu", argv[i])==0){
nrgpu = atoi(argv[++i]);
if (nrgpu>0){
printf(" testing_zhegvdx -ngpu %d\n\n", (int) nrgpu);
}
else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -ngpu %d\n\n", (int) nrgpu);
exit(1);
}
}
if (strcmp("-itype", argv[i])==0){
itype = atoi(argv[++i]);
if (itype>0 && itype <= 3){
printf(" testing_zhegvdx -itype %d\n\n", (int) itype);
}
else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -itype %d\n\n", (int) itype);
exit(1);
}
}
if (strcmp("-FE", argv[i])==0){
fraction_ev = atof(argv[++i]);
if (fraction_ev > 0 && fraction_ev <= 1){
printf(" testing_zhegvdx -FE %f\n\n", fraction_ev);
}
else {
fraction_ev = 0;
}
}
if (strcmp("-L", argv[i])==0){
uplo = (char*)MagmaLowerStr;
printf(" testing_zhegvdx -L");
}
if (strcmp("-U", argv[i])==0){
uplo = (char*)MagmaUpperStr;
printf(" testing_zhegvdx -U");
}
}
} else {
printf("\nUsage: \n");
printf(" testing_zhegvdx -L/U -N %d -itype %d\n\n", 1024, 1);
}
if(!flagN)
N = size[3];
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
n2 = N * N;
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, cuDoubleComplex, n2);
TESTING_MALLOC( h_B, cuDoubleComplex, n2);
TESTING_MALLOC( w1, double , N);
TESTING_MALLOC( w2, double , N);
TESTING_HOSTALLOC(h_R, cuDoubleComplex, n2);
TESTING_HOSTALLOC(h_S, cuDoubleComplex, n2);
magma_int_t nb = magma_get_zhetrd_nb(N);
magma_int_t lwork = magma_zbulge_get_lq2(N) + 2*N + N*N;
magma_int_t lrwork = 1 + 5*N +2*N*N;
magma_int_t liwork = 3 + 5*N;
TESTING_HOSTALLOC(h_work, cuDoubleComplex, lwork);
TESTING_HOSTALLOC( rwork, double, lrwork);
TESTING_MALLOC( iwork, magma_int_t, liwork);
printf(" N M GPU Time(s) \n");
printf("==========================\n");
for(i=0; i<4; i++){
if (!flagN){
N = size[i];
n2 = N*N;
}
if (fraction_ev == 0){
il = N / 10;
iu = N / 5+il;
}
else {
il = 1;
iu = (int)(fraction_ev*N);
if (iu < 1) iu = 1;
}
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
//lapackf77_zlatms( &N, &N, "U", ISEED, "P", w1, &five, &d_ten,
// &d_one, &N, &N, uplo, h_B, &N, h_work, &info);
//lapackf77_zlaset( "A", &N, &N, &c_zero, &c_one, h_B, &N);
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_Z_SET2REAL( h_B[i*N+i], ( MAGMA_Z_REAL(h_B[i*N+i]) + 1.*N ) );
}
}
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
#ifdef USE_MGPU
magma_zhegvdx_2stage_m(nrgpu, itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#else
magma_zhegvdx_2stage(itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#endif
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
#ifdef USE_MGPU
magma_zhegvdx_2stage_m(nrgpu, itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#else
magma_zhegvdx_2stage(itype, jobz[0], range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
#endif
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | S(with V) - S(w/o V) | / | S |
=================================================================== */
double temp1, temp2;
cuDoubleComplex *tau;
result[0] = 1.;
result[0] /= lapackf77_zlanhe("1",uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_zlange("1",&N , &m1, h_R, &N, rwork);
if (itype == 1){
blasf77_zhemm("L", uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", uplo, &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (itype == 2){
blasf77_zhemm("L", uplo, &N, &m1, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", uplo, &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (itype == 3){
blasf77_zhemm("L", uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", uplo, &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_zlange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_zhegvdx(itype, 'N', range, uplo[0],
N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
temp1 = temp2 = 0;
for(int j=0; j<m2; j++){
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[1] = temp2 / temp1;
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %6.2f\n",
(int) N, (int) m1, gpu_time);
if ( checkres ){
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
printf("(2) | D(w/ Z)-D(w/o Z)|/ |D| = %e\n\n", result[1]);
}
if (flagN)
break;
}
cudaSetDevice(0);
/* Memory clean up */
TESTING_FREE( h_A);
TESTING_FREE( h_B);
TESTING_FREE( w1);
TESTING_FREE( w2);
TESTING_HOSTFREE( rwork);
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
TESTING_HOSTFREE( h_S);
/* Shutdown */
#ifdef USE_MGPU
TESTING_CUDA_FINALIZE_MGPU();
#else
TESTING_CUDA_FINALIZE();
#endif
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cpotrf
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
float flops, gpu_perf, cpu_perf;
cuFloatComplex *h_A, *h_R;
cuFloatComplex *d_A;
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info;
const char *uplo = MagmaUpperStr;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float work[1], matnorm;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
}
if (N>0) size[0] = size[9] = N;
else exit(1);
}
else {
printf("\nUsage: \n");
printf(" testing_cpotri_gpu -N %d\n\n", 1024);
}
/* Allocate host memory for the matrix */
n2 = size[9] * size[9];
ldda = ((size[9]+31)/32) * 32;
TESTING_MALLOC( h_A, cuFloatComplex, n2);
TESTING_HOSTALLOC( h_R, cuFloatComplex, n2);
TESTING_DEVALLOC( d_A, cuFloatComplex, ldda*size[9] );
printf(" N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("========================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
flops = FLOPS_CPOTRI( (float)N ) / 1000000;
ldda = ((N+31)/32)*32;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
/* Symmetrize and increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_A[i*lda+i], ( MAGMA_C_REAL(h_A[i*lda+i]) + 1.*N ) );
for(j=0; j<i; j++)
h_A[i*lda+j] = cuConjf(h_A[j*lda+i]);
}
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
//cublasSetMatrix( N, N, sizeof(cuFloatComplex), h_A, lda, d_A, ldda);
//magma_cpotrf_gpu(uplo[0], N, d_A, ldda, &info);
/* factorize matrix */
magma_csetmatrix( N, N, h_A, lda, d_A, ldda );
magma_cpotrf_gpu(uplo[0], N, d_A, ldda, &info);
// check for exact singularity
//magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
//h_R[ 10 + 10*lda ] = MAGMA_C_MAKE( 0.0, 0.0 );
//magma_csetmatrix( N, N, h_R, lda, d_A, ldda );
start = get_current_time();
magma_cpotri_gpu(uplo[0], N, d_A, ldda, &info);
end = get_current_time();
if (info != 0)
printf("magma_cpotri_gpu returned error %d\n", (int) info);
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_cpotrf(uplo, &N, h_A, &lda, &info);
start = get_current_time();
lapackf77_cpotri(uplo, &N, h_A, &lda, &info);
end = get_current_time();
if (info != 0)
printf("lapackf77_cpotri returned error %d\n", (int) info);
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( N, N, d_A, ldda, h_R, lda );
matnorm = lapackf77_clange("f", &N, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %6.2f %6.2f %e\n",
(int) size[i], cpu_perf, gpu_perf,
lapackf77_clange("f", &N, &N, h_R, &lda, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magma_int_t M, N, n2, lda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol, eps = lapackf77_dlamch("E");
tol = opts.tolerance * eps;
opts.lapack |= ( opts.check == 2 ); // check (-c2) implies lapack (-l)
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("===============================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
}
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
nb = magma_get_zgeqrf_nb(M);
gflops = FLOPS_ZGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_zgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, max( N*nb, 2*nb*nb ));
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC( h_work, magmaDoubleComplex, lwork );
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgeqrf(M, N, h_R, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
magmaDoubleComplex *tau;
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
cpu_time = magma_wtime();
lapackf77_zgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE( tau );
}
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
magmaDoubleComplex *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC( h_W1, magmaDoubleComplex, n2 ); // Q
TESTING_MALLOC( h_W2, magmaDoubleComplex, n2 ); // R
TESTING_MALLOC( h_W3, magmaDoubleComplex, lwork ); // WORK
TESTING_MALLOC( h_RW, double, M ); // RWORK
lapackf77_zlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_zqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
}
status |= ! (results[0] < tol);
TESTING_FREE( h_W1 );
TESTING_FREE( h_W2 );
TESTING_FREE( h_W3 );
TESTING_FREE( h_RW );
} else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &M, &N, h_R, &lda, work) / error;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf("%s\n", (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf_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];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_R, *tau, *h_work, tmp[1];
double *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 i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_dgeqrf_nb( M );
gflops = FLOPS_DGEQRF( 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_dgeqrf( &M, &N, h_A, &M, tau, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_D_REAL( tmp[0] );
// Allocate host memory for the matrix
TESTING_MALLOC( tau, double, min_mn );
TESTING_MALLOC( h_A, double, n2 );
TESTING_HOSTALLOC( h_R, double, n2 );
TESTING_MALLOC( h_work, double, lhwork );
// 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_DEVALLOC( d_lA[dev], double, ldda*n_local );
}
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
double *tau;
TESTING_MALLOC( tau, double, min_mn );
cpu_time = magma_wtime();
lapackf77_dgeqrf( &M, &N, h_A, &M, tau, h_work, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE( tau );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix_1D_col_bcyclic( M, N, h_R, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_dgeqrf2_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_dgeqrf2 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_R, lda, ngpu, nb );
magma_queue_sync( NULL );
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
double *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC( h_W1, double, n2 ); // Q
TESTING_MALLOC( h_W2, double, n2 ); // R
TESTING_MALLOC( h_W3, double, lwork ); // WORK
TESTING_MALLOC( h_RW, double, M ); // RWORK
lapackf77_dlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_dqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
}
status |= ! (results[0] < tol);
TESTING_FREE( h_W1 );
TESTING_FREE( h_W2 );
TESTING_FREE( h_W3 );
TESTING_FREE( h_RW );
} else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work );
blasf77_daxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work ) / (min_mn*error);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
printf("%s\n", (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_DEVFREE( d_lA[dev] );
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error;
magmaDoubleComplex *h_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |Ax-b|/(N*|A|*|x|)\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |PA-LU|/(N*|A|)\n");
}
printf("=========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_ZGETRF( M, N ) / 1e9;
TESTING_MALLOC( ipiv, magma_int_t, min_mn );
TESTING_HOSTALLOC( h_A, magmaDoubleComplex, n2 );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_zgetrf(&M, &N, h_A, &lda, ipiv, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( M, N, h_A, lda );
gpu_time = magma_wtime();
magma_zgetrf( M, N, h_A, lda, ipiv, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time );
}
if ( opts.check == 2 ) {
error = get_residual( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else if ( opts.check ) {
error = get_LU_error( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
printf(" --- \n");
}
TESTING_FREE( ipiv );
TESTING_HOSTFREE( h_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqlf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_zgeqrf_nb(M);
gflops = FLOPS_ZGEQLF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_zgeqlf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max( lwork, N*nb );
TESTING_MALLOC( tau, magmaDoubleComplex, min_mn );
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC( h_work, magmaDoubleComplex, lwork );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgeqlf( M, N, h_R, lda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zgeqlf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_zgeqlf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_zlange("f", &M, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_FREE( h_work );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
magmaFloatComplex *d_A;
float *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9;
lwork = -1;
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC( h_A, magmaFloatComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaFloatComplex, n2 );
TESTING_DEVALLOC( d_A, magmaFloatComplex, ldda*N );
TESTING_DEVALLOC( dtau, magmaFloatComplex, min_mn );
TESTING_DEVALLOC(dwork, float, min_mn );
TESTING_MALLOC( h_work, magmaFloatComplex, lwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_cgeqr2_gpu( M, N, d_A, ldda, dtau, dwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
error = lapackf77_clange("f", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_clange("f", &M, &N, h_R, &lda, work) / error;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, error );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( dtau );
TESTING_DEVFREE( dwork );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgetrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error;
float *h_A, *h_R;
float *d_A;
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||PA-LU||/(||A||*N)\n");
printf("=========================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_SGETRF( M, N ) / 1e9;
TESTING_MALLOC( ipiv, magma_int_t, min_mn );
TESTING_MALLOC( h_A, float, n2 );
TESTING_HOSTALLOC( h_R, float, n2 );
TESTING_DEVALLOC( d_A, float, ldda*N );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_sgetrf(&M, &N, h_A, &lda, ipiv, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_sgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_sgetf2_gpu( M, N, d_A, ldda, ipiv, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_sgetf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000. );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time*1000. );
}
if ( opts.check ) {
magma_sgetmatrix( M, N, d_A, ldda, h_A, lda );
error = get_LU_error( M, N, h_R, lda, h_A, ipiv );
printf(" %8.2e\n", error );
}
else {
printf(" --- \n");
}
TESTING_FREE( ipiv );
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dswap, dswapblk, dpermute, dlaswp, dlaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
double *h_A1, *h_A2;
double *d_A1, *d_A2;
double *h_R1, *h_R2;
// row-major and column-major performance
real_Double_t row_perf0, col_perf0;
real_Double_t row_perf1, col_perf1;
real_Double_t row_perf2, col_perf2;
real_Double_t row_perf3;
real_Double_t row_perf4;
real_Double_t row_perf5, col_perf5;
real_Double_t row_perf6, col_perf6;
real_Double_t row_perf7;
real_Double_t cpu_perf;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magma_int_t *d_ipiv;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasDswap dswap dswapblk dlaswp dpermute dlaswp2 dlaswpx dcopymatrix CPU (all in )\n");
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj row-maj/col-maj row-blk/col-blk dlaswp (GByte/s)\n");
printf("==================================================================================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// each test is assigned one bit in the check bitmask, bit=1 is failure.
// shift keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[i];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_dgetrf_nb( N ));
// for each swap, does 2N loads and 2N stores
gbytes = sizeof(double) * 4.*N*nb / 1e9;
TESTING_HOSTALLOC( h_A1, double, lda*N );
TESTING_HOSTALLOC( h_A2, double, lda*N );
TESTING_HOSTALLOC( h_R1, double, lda*N );
TESTING_HOSTALLOC( h_R2, double, lda*N );
TESTING_MALLOC( ipiv, magma_int_t, nb );
TESTING_MALLOC( ipiv2, magma_int_t, nb );
TESTING_DEVALLOC( d_ipiv, magma_int_t, nb );
TESTING_DEVALLOC( d_A1, double, ldda*N );
TESTING_DEVALLOC( d_A2, double, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasDswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasDswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasDswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda);
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* dswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_dswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_dswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* dswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_dswapblk( 'R', N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_dswapblk( 'C', N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* dpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // dpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dpermute_long2( N, d_A1, ldda, ipiv2, nb, 0 );
time = magma_sync_wtime( queue ) - time;
row_perf3 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_dlaswp2( N, d_A1, ldda, 1, nb, d_ipiv );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_dlaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_dcopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_dcopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf3, ((check & 0x040) != 0 ? '*' : ' '),
row_perf4, ((check & 0x080) != 0 ? '*' : ' '),
row_perf7, ((check & 0x100) != 0 ? '*' : ' '),
row_perf5, ((check & 0x200) != 0 ? '*' : ' '),
col_perf5, ((check & 0x400) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failures") );
TESTING_HOSTFREE( h_A1 );
TESTING_HOSTFREE( h_A2 );
TESTING_HOSTFREE( h_R1 );
TESTING_HOSTFREE( h_R2 );
TESTING_DEVFREE( d_A1 );
TESTING_DEVFREE( d_A2 );
TESTING_FREE( ipiv );
TESTING_FREE( ipiv2 );
TESTING_DEVFREE( d_ipiv );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cungqr_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *hA, *hR, *tau, *h_work;
magmaFloatComplex *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_cgeqrf_nb( m );
lwork = (m + 2*n+nb)*nb;
gflops = FLOPS_CUNGQR( m, n, k ) / 1e9;
TESTING_HOSTALLOC( hA, magmaFloatComplex, lda*n );
TESTING_HOSTALLOC( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC( hR, magmaFloatComplex, lda*n );
TESTING_MALLOC( tau, magmaFloatComplex, min_mn );
TESTING_DEVALLOC( dA, magmaFloatComplex, ldda*n );
TESTING_DEVALLOC( dT, magmaFloatComplex, ( 2*min_mn + ((n + 31)/32)*32 )*nb );
lapackf77_clarnv( &ione, ISEED, &n2, hA );
lapackf77_clacpy( MagmaUpperLowerStr, &m, &n, hA, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( m, n, hA, lda, dA, ldda );
magma_cgeqrf_gpu( m, n, dA, ldda, tau, dT, &info );
if (info != 0)
printf("magma_cgeqrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
gpu_time = magma_wtime();
magma_cungqr_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_cungqr_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get dA back to the CPU to compare with the CPU result.
magma_cgetmatrix( m, n, dA, ldda, hR, lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
error = lapackf77_clange("f", &m, &n, hA, &lda, work );
lapackf77_cgeqrf( &m, &n, hA, &lda, tau, h_work, &lwork, &info );
if (info != 0)
printf("lapackf77_cgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
cpu_time = magma_wtime();
lapackf77_cungqr( &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_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute relative error |R|/|A| := |Q_magma - Q_lapack|/|A|
blasf77_caxpy( &n2, &c_neg_one, hA, &ione, hR, &ione );
error = lapackf77_clange("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_HOSTFREE( hA );
TESTING_HOSTFREE( h_work );
TESTING_FREE( hR );
TESTING_FREE( tau );
TESTING_DEVFREE( dA );
TESTING_DEVFREE( dT );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time;
magmaFloatComplex *h_A, *h_R, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork;
#endif
/* Matrix size */
float *w1, *w2;
magma_int_t *iwork;
magma_int_t N, n2, info, lwork, liwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};;
magma_int_t info_ortho = 0;
magma_int_t info_solution = 0;
magma_int_t info_reduction = 0;
magma_timestr_t start, end;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t ngpu = opts.ngpu;
char jobz = opts.jobz;
magma_int_t checkres = opts.check;
char range = 'A';
char uplo = opts.uplo;
magma_int_t itype = opts.itype;
float f = opts.fraction;
if (f != 1)
range='I';
if ( checkres && jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
jobz = MagmaVec;
}
printf("using: itype = %d, jobz = %c, range = %c, uplo = %c, checkres = %d, fraction = %6.4f\n",
(int) itype, jobz, range, uplo, (int) checkres, f);
printf(" N M GPU Time(s) \n");
printf("==========================\n");
magma_int_t threads = magma_get_numthreads();
for( magma_int_t i = 0; i < opts.ntest; ++i ) {
for( magma_int_t iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = magma_cbulge_get_lq2(N, threads) + 2*N + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = magma_cbulge_get_lq2(N, threads) + 1 + 6*N + 2*N*N;
#endif
liwork = 3 + 5*N;
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, magmaFloatComplex, n2);
TESTING_MALLOC( w1, float , N);
TESTING_MALLOC( w2, float , N);
TESTING_HOSTALLOC(h_R, magmaFloatComplex, n2);
TESTING_HOSTALLOC(h_work, magmaFloatComplex, lwork);
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTALLOC( rwork, float, lrwork);
#endif
TESTING_MALLOC( iwork, magma_int_t, liwork);
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
/* Make diagonal real */
for(int i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_A[i*N+i], MAGMA_C_REAL(h_A[i*N+i]) );
}
magma_int_t m1 = 0;
float vl = 0;
float vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (range == 'I'){
il = 1;
iu = (int) (f*N);
}
if(opts.warmup){
// ==================================================================
// Warmup using MAGMA
// ==================================================================
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
if(ngpu==1){
printf("calling cheevdx_2stage 1 GPU\n");
magma_cheevdx_2stage(jobz, range, uplo, N,
h_R, N,
vl, vu, il, iu,
&m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
}else{
printf("calling cheevdx_2stage_m %d GPU\n", (int) ngpu);
magma_cheevdx_2stage_m(ngpu, jobz, range, uplo, N,
h_R, N,
vl, vu, il, iu,
&m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
}
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
start = get_current_time();
if(ngpu==1){
printf("calling cheevdx_2stage 1 GPU\n");
magma_cheevdx_2stage(jobz, range, uplo, N,
h_R, N,
vl, vu, il, iu,
&m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
}else{
printf("calling cheevdx_2stage_m %d GPU\n", (int) ngpu);
magma_cheevdx_2stage_m(ngpu, jobz, range, uplo, N,
h_R, N,
vl, vu, il, iu,
&m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
}
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
float eps = lapackf77_slamch("E");
printf("\n");
printf("------ TESTS FOR MAGMA CHEEVD ROUTINE ------- \n");
printf(" Size of the Matrix %d by %d\n", (int) N, (int) N);
printf("\n");
printf(" The matrix A is randomly generated for each test.\n");
printf("============\n");
printf(" The relative machine precision (eps) is %8.2e\n",eps);
printf(" Computational tests pass if scaled residuals are less than 60.\n");
/* Check the orthogonality, reduction and the eigen solutions */
if (jobz == MagmaVec) {
info_ortho = check_orthogonality(N, N, h_R, N, eps);
info_reduction = check_reduction(uplo, N, 1, h_A, w1, N, h_R, eps);
}
printf("------ CALLING LAPACK CHEEVD TO COMPUTE only eigenvalue and verify elementswise ------- \n");
lapackf77_cheevd("N", "L", &N,
h_A, &N, w2,
h_work, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork,
&info);
info_solution = check_solution(N, w2, w1, eps);
if ( (info_solution == 0) & (info_ortho == 0) & (info_reduction == 0) ) {
printf("***************************************************\n");
printf(" ---- TESTING CHEEVD ...................... PASSED !\n");
printf("***************************************************\n");
}
else {
printf("************************************************\n");
printf(" - TESTING CHEEVD ... FAILED !\n");
printf("************************************************\n");
}
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %6.2f\n",
(int) N, (int) m1, gpu_time);
TESTING_FREE( h_A);
TESTING_FREE( w1);
TESTING_FREE( w2);
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTFREE( rwork);
#endif
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgehrd
*/
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, *h_Q, *h_work, *tau, *twork, *dT;
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
#endif
double eps, result[2];
magma_int_t N, n2, lda, nb, lwork, ltwork, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
eps = lapackf77_dlamch( "E" );
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) |A-QHQ'|/N|A| |I-QQ'|/N\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;
nb = magma_get_zgehrd_nb(N);
/* We suppose the magma nb is bigger than lapack nb */
lwork = N*nb;
gflops = FLOPS_ZGEHRD( N ) / 1e9;
TESTING_MALLOC ( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC ( tau, magmaDoubleComplex, N );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_work, magmaDoubleComplex, lwork );
TESTING_DEVALLOC ( dT, magmaDoubleComplex, nb*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zgehrd( N, ione, N, h_R, lda, tau, h_work, lwork, dT, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zgehrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.check ) {
ltwork = 2*(N*N);
TESTING_HOSTALLOC( h_Q, magmaDoubleComplex, lda*N );
TESTING_MALLOC( twork, magmaDoubleComplex, ltwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC( rwork, double, N );
#endif
lapackf77_zlacpy(MagmaUpperLowerStr, &N, &N, h_R, &lda, h_Q, &lda);
for( int j = 0; j < N-1; ++j )
for( int i = j+2; i < N; ++i )
h_R[i+j*lda] = MAGMA_Z_ZERO;
magma_zunghr(N, ione, N, h_Q, lda, tau, dT, nb, &info);
if (info != 0) {
printf("magma_zunghr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
exit(1);
}
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zhst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, rwork, result);
#else
lapackf77_zhst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, result);
#endif
TESTING_HOSTFREE( h_Q );
TESTING_FREE( twork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE( rwork );
#endif
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zgehrd(&N, &ione, &N, h_R, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zgehrd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Print performance and error.
=================================================================== */
if ( opts.lapack ) {
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f)",
(int) N, gpu_perf, gpu_time );
}
if ( opts.check ) {
printf(" %8.2e %8.2e%s\n",
result[0]*eps, result[1]*eps,
( ( (result[0]*eps < tol) && (result[1]*eps < tol) ) ? "" : " failed") );
status |= ! (result[0]*eps < tol);
status |= ! (result[1]*eps < tol);
}
else {
printf(" --- ---\n");
}
TESTING_FREE ( h_A );
TESTING_FREE ( tau );
TESTING_HOSTFREE( h_work);
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE ( dT );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvd
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time, cpu_time;
magmaDoubleComplex *h_A, *h_R, *h_B, *h_S, *h_work;
double *rwork, *w1, *w2;
double result[4] = {0};
magma_int_t *iwork;
magma_int_t N, n2, info, nb, lwork, liwork, lda, lrwork;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double d_one = 1.;
double d_neg_one = -1.;
//double d_ten = 10.;
//magma_int_t izero = 0;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
double tolulp = opts.tolerance * lapackf77_dlamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf(" N CPU Time (sec) GPU Time(sec)\n");
printf("======================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = N*lda;
nb = magma_get_zhetrd_nb(N);
lwork = 2*N*nb + N*N;
lrwork = 1 + 5*N +2*N*N;
liwork = 3 + 5*N;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC( h_B, magmaDoubleComplex, n2 );
TESTING_MALLOC( w1, double, N );
TESTING_MALLOC( w2, double, N );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_S, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_work, magmaDoubleComplex, lwork );
TESTING_MALLOC( rwork, double, lrwork );
TESTING_MALLOC( iwork, magma_int_t, liwork );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
//lapackf77_zlatms( &N, &N, "U", ISEED, "P", w1, &five, &d_ten,
// &d_one, &N, &N, &opts.uplo, h_B, &lda, h_work, &info);
//lapackf77_zlaset( "A", &N, &N, &c_zero, &c_one, h_B, &lda);
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
magma_zmake_hpd( N, h_B, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda );
/* warmup */
if ( opts.warmup ) {
magma_zhegvd( opts.itype, opts.jobz, opts.uplo,
N, h_R, lda, h_S, lda, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info );
if (info != 0)
printf("magma_zhegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zhegvd( opts.itype, opts.jobz, opts.uplo,
N, h_R, lda, h_S, lda, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_zhegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvd routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | I - V V' B | / ( N ) (itype = 1,2)
| B - V V' | / ( |B| N ) (itype = 3)
(3) | S(with V) - S(w/o V) | / | S |
=================================================================== */
double temp1, temp2;
//magmaDoubleComplex *tau;
if ( opts.itype == 1 || opts.itype == 2 ) {
lapackf77_zlaset( "A", &N, &N, &c_zero, &c_one, h_S, &lda);
blasf77_zgemm("N", "C", &N, &N, &N, &c_one, h_R, &lda, h_R, &lda, &c_zero, h_work, &N);
blasf77_zhemm("R", &opts.uplo, &N, &N, &c_neg_one, h_B, &lda, h_work, &N, &c_one, h_S, &lda);
result[1] = lapackf77_zlange("1", &N, &N, h_S, &lda, rwork) / N;
}
else if ( opts.itype == 3 ) {
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda);
blasf77_zherk(&opts.uplo, "N", &N, &N, &d_neg_one, h_R, &lda, &d_one, h_S, &lda);
result[1] = lapackf77_zlanhe("1", &opts.uplo, &N, h_S, &lda, rwork) / N
/ lapackf77_zlanhe("1", &opts.uplo, &N, h_B, &lda, rwork);
}
result[0] = 1.;
result[0] /= lapackf77_zlanhe("1", &opts.uplo, &N, h_A, &lda, rwork);
result[0] /= lapackf77_zlange("1", &N, &N, h_R, &lda, rwork);
if ( opts.itype == 1 ) {
blasf77_zhemm("L", &opts.uplo, &N, &N, &c_one, h_A, &lda, h_R, &lda, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", &opts.uplo, &N, &N, &c_neg_one, h_B, &lda, h_R, &lda, &c_one, h_work, &N);
result[0] *= lapackf77_zlange("1", &N, &N, h_work, &lda, rwork)/N;
}
else if ( opts.itype == 2 ) {
blasf77_zhemm("L", &opts.uplo, &N, &N, &c_one, h_B, &lda, h_R, &lda, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", &opts.uplo, &N, &N, &c_one, h_A, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_zlange("1", &N, &N, h_R, &lda, rwork)/N;
}
else if ( opts.itype == 3 ) {
blasf77_zhemm("L", &opts.uplo, &N, &N, &c_one, h_A, &lda, h_R, &lda, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_zdscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_zhemm("L", &opts.uplo, &N, &N, &c_one, h_B, &lda, h_work, &N, &c_neg_one, h_R, &lda);
result[0] *= lapackf77_zlange("1", &N, &N, h_R, &lda, rwork)/N;
}
/*
lapackf77_zhet21( &ione, &opts.uplo, &N, &izero,
h_A, &lda,
w1, w1,
h_R, &lda,
h_R, &lda,
tau, h_work, rwork, &result[0] );
*/
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &lda, h_S, &lda );
magma_zhegvd( opts.itype, MagmaNoVec, opts.uplo,
N, h_R, lda, h_S, lda, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info );
if (info != 0)
printf("magma_zhegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for(int j=0; j<N; j++) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / (((double)N)*temp1);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_zhegvd( &opts.itype, &opts.jobz, &opts.uplo,
&N, h_A, &lda, h_B, &lda, w2,
h_work, &lwork,
rwork, &lrwork,
iwork, &liwork,
&info );
cpu_time = magma_wtime() - cpu_time;
if (info != 0)
printf("lapackf77_zhegvd returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf("%5d %7.2f %7.2f\n",
(int) N, cpu_time, gpu_time);
}
else {
printf("%5d --- %7.2f\n",
(int) N, gpu_time);
}
/* =====================================================================
Print execution time
=================================================================== */
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if ( opts.itype==1 )
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed") );
else if ( opts.itype==2 )
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed") );
else if ( opts.itype==3 )
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed") );
if ( opts.itype==1 || opts.itype==2 )
printf("(2) | I - Z Z' B | / N = %8.2e%s\n", result[1], (result[1] < tol ? "" : " failed") );
else
printf("(2) | B - Z Z' | / (|B| N) = %8.2e%s\n", result[1], (result[1] < tol ? "" : " failed") );
printf( "(3) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e%s\n\n", result[2], (result[2] < tolulp ? "" : " failed") );
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( w1 );
TESTING_FREE( w2 );
TESTING_FREE( rwork );
TESTING_FREE( iwork );
TESTING_HOSTFREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_HOSTFREE( h_S );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgehrd2
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
double eps, flops, gpu_perf, cpu_perf;
cuDoubleComplex *h_A, *h_R, *h_Q, *h_work, *tau, *twork, *dT;
double *rwork;
double result[2] = {0., 0.};
/* Matrix size */
magma_int_t N=0, n2, lda, nb, lwork, ltwork, once = 0;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info, checkres;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0) {
N = atoi(argv[++i]);
once = true;
}
}
if ( N > 0 )
printf(" testing_zgehrd -N %d\n\n", (int) N);
else
{
printf("\nUsage: \n");
printf(" testing_zgehrd -N %d\n\n", 1024);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_zgehrd -N %d\n\n", 1024);
N = size[9];
}
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
eps = lapackf77_dlamch( "E" );
lda = N;
n2 = N*lda;
nb = magma_get_zgehrd_nb(N);
/* We suppose the magma nb is bigger than lapack nb */
lwork = N*nb;
TESTING_MALLOC ( h_A , cuDoubleComplex, n2 );
TESTING_MALLOC ( tau , cuDoubleComplex, N );
TESTING_HOSTALLOC( h_R , cuDoubleComplex, n2 );
TESTING_HOSTALLOC( h_work, cuDoubleComplex, lwork );
TESTING_DEVALLOC ( dT , cuDoubleComplex, nb*N );
/* To avoid uninitialized variable warning */
h_Q = NULL;
twork = NULL;
rwork = NULL;
if ( checkres ) {
ltwork = 2*(N*N);
TESTING_HOSTALLOC( h_Q, cuDoubleComplex, lda*N );
TESTING_MALLOC( twork, cuDoubleComplex, ltwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_MALLOC( rwork, double, N );
#endif
}
printf(" N CPU GFlop/s GPU GFlop/s |A-QHQ'|/N|A| |I-QQ'|/N \n");
printf("=============================================================\n");
for(i=0; i<10; i++){
if ( !once ) {
N = size[i];
}
lda = N;
n2 = lda*N;
flops = FLOPS( (double)N ) / 1e6;
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_zgehrd ( N, ione, N, h_R, lda, tau, h_work, lwork, dT, &info);
end = get_current_time();
if ( info < 0 )
printf("Argument %d of magma_zgehrd had an illegal value\n", (int) -info);
gpu_perf = flops / GetTimerValue(start,end);
/* =====================================================================
Check the factorization
=================================================================== */
if ( checkres ) {
lapackf77_zlacpy(MagmaUpperLowerStr, &N, &N, h_R, &lda, h_Q, &lda);
{
int i, j;
for(j=0; j<N-1; j++)
for(i=j+2; i<lda; i++)
h_R[i+j*lda] = MAGMA_Z_ZERO;
}
nb = magma_get_zgehrd_nb(N);
magma_zunghr(N, ione, N, h_Q, lda, tau, dT, nb, &info);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zhst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, rwork, result);
#else
lapackf77_zhst01(&N, &ione, &N,
h_A, &lda, h_R, &lda,
h_Q, &lda, twork, <work, result);
#endif
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_zgehrd(&N, &ione, &N, h_R, &lda, tau, h_work, &lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of lapack_zgehrd had an illegal value.\n", (int) -info);
cpu_perf = flops / GetTimerValue(start,end);
/* =====================================================================
Print performance and error.
=================================================================== */
if ( checkres ) {
printf("%5d %6.2f %6.2f %e %e\n",
(int) N, cpu_perf, gpu_perf,
result[0]*eps, result[1]*eps );
} else {
printf("%5d %6.2f %6.2f\n",
(int) N, cpu_perf, gpu_perf );
}
if ( once )
break;
}
/* Memory clean up */
TESTING_FREE ( h_A );
TESTING_FREE ( tau );
TESTING_HOSTFREE( h_work);
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE ( dT );
if ( checkres ) {
TESTING_HOSTFREE( h_Q );
TESTING_FREE( twork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_FREE( rwork );
#endif
}
/* Shutdown */
TESTING_CUDA_FINALIZE();
return EXIT_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zhegvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT_MGPU();
real_Double_t mgpu_time;
magmaDoubleComplex *h_A, *h_Ainit, *h_B, *h_Binit, *h_work;
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork;
magma_int_t lrwork;
#endif
double *w1, result;
magma_int_t *iwork;
magma_int_t N, n2, info, lwork, liwork;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_timestr_t start, end;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
char jobz = opts.jobz;
int checkres = opts.check;
char range = 'A';
char uplo = opts.uplo;
magma_int_t itype = opts.itype;
double f = opts.fraction;
if (f != 1)
range='I';
if ( checkres && jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
jobz = MagmaVec;
}
printf("using: nrgpu = %d, itype = %d, jobz = %c, range = %c, uplo = %c, checkres = %d, fraction = %6.4f\n",
(int) opts.ngpu, (int) itype, jobz, range, uplo, (int) checkres, f);
printf(" N M nr GPU MGPU Time(s) \n");
printf("====================================\n");
magma_int_t threads = magma_get_numthreads();
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = magma_zbulge_get_lq2(N, threads) + 2*N + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = magma_zbulge_get_lq2(N, threads) + 1 + 6*N + 2*N*N;
#endif
liwork = 3 + 5*N;
//magma_int_t NB = 96;//magma_bulge_get_nb(N);
//magma_int_t sizvblg = magma_zbulge_get_lq2(N, threads);
//magma_int_t siz = max(sizvblg,n2)+2*(N*NB+N)+24*N;
/* Allocate host memory for the matrix */
TESTING_HOSTALLOC( h_A, magmaDoubleComplex, n2);
TESTING_HOSTALLOC( h_B, magmaDoubleComplex, n2);
TESTING_MALLOC( w1, double , N);
TESTING_HOSTALLOC(h_work, magmaDoubleComplex, lwork);
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTALLOC( rwork, double, lrwork);
#endif
TESTING_MALLOC( iwork, magma_int_t, liwork);
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
{
for(int i=0; i<N; i++) {
MAGMA_Z_SET2REAL( h_B[i*N+i], ( MAGMA_Z_REAL(h_B[i*N+i]) + 1.*N ) );
MAGMA_Z_SET2REAL( h_A[i*N+i], MAGMA_Z_REAL(h_A[i*N+i]) );
}
}
if((opts.warmup)||( checkres )){
TESTING_MALLOC(h_Ainit, magmaDoubleComplex, n2);
TESTING_MALLOC(h_Binit, magmaDoubleComplex, n2);
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_Ainit, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_Binit, &N );
}
magma_int_t m1 = 0;
double vl = 0;
double vu = 0;
magma_int_t il = 0;
magma_int_t iu = 0;
if (range == 'I'){
il = 1;
iu = (int) (f*N);
}
if(opts.warmup){
// ==================================================================
// Warmup using MAGMA. I prefer to use smalltest to warmup A-
// ==================================================================
magma_zhegvdx_2stage_m(opts.ngpu, itype, jobz, range, uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_Ainit, &N, h_A, &N );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_Binit, &N, h_B, &N );
}
// ===================================================================
// Performs operation using MAGMA
// ===================================================================
start = get_current_time();
magma_zhegvdx_2stage_m(opts.ngpu, itype, jobz, range, uplo,
N, h_A, N, h_B, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info);
end = get_current_time();
mgpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
// ===================================================================
// Check the results following the LAPACK's [zc]hegvdx routine.
// A x = lambda B x is solved
// and the following 3 tests computed:
// (1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
// | A B Z - Z D | / ( |A||Z| N ) (itype = 2)
// | B A Z - Z D | / ( |A||Z| N ) (itype = 3)
// ===================================================================
#if defined(PRECISION_d) || defined(PRECISION_s)
double *rwork = h_work + N*N;
#endif
result = 1.;
result /= lapackf77_zlanhe("1",&uplo, &N, h_Ainit, &N, rwork);
result /= lapackf77_zlange("1",&N , &m1, h_A, &N, rwork);
if (itype == 1){
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_zhemm("L", &uplo, &N, &m1, &c_neg_one, h_Binit, &N, h_A, &N, &c_one, h_work, &N);
result *= lapackf77_zlange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (itype == 2){
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Binit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Ainit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_zlange("1", &N, &m1, h_A, &N, rwork)/N;
}
else if (itype == 3){
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Ainit, &N, h_A, &N, &c_zero, h_work, &N);
for(int i=0; i<m1; ++i)
blasf77_zdscal(&N, &w1[i], &h_A[i*N], &ione);
blasf77_zhemm("L", &uplo, &N, &m1, &c_one, h_Binit, &N, h_work, &N, &c_neg_one, h_A, &N);
result *= lapackf77_zlange("1", &N, &m1, h_A, &N, rwork)/N;
}
}
// ===================================================================
// Print execution time
// ===================================================================
printf("%5d %5d %2d %6.2f\n",
(int) N, (int) m1, (int) opts.ngpu, mgpu_time);
if ( checkres ){
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result, (result < tol ? "" : " failed") );
}
TESTING_HOSTFREE( h_A);
TESTING_HOSTFREE( h_B);
TESTING_FREE( w1);
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTFREE( rwork);
#endif
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
if((opts.warmup)||( checkres )){
TESTING_FREE( h_Ainit);
TESTING_FREE( h_Binit);
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
/* Shutdown */
TESTING_FINALIZE_MGPU();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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 cheevd
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
cuFloatComplex *h_A, *h_R, *h_work;
float *rwork, *w1, *w2;
magma_int_t *iwork;
float gpu_time, cpu_time;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t size[8] = {1024,2048,3072,4032,5184,6016,7040,8064};
magma_int_t i, info;
magma_int_t ione = 1, izero = 0;
magma_int_t ISEED[4] = {0,0,0,1};
const char *uplo = MagmaLowerStr;
const char *jobz = MagmaVectorsStr;
magma_int_t checkres;
float result[3], eps = lapackf77_slamch( "E" );
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0) {
N = atoi(argv[++i]);
}
else if ( strcmp("-JV", argv[i]) == 0 ) {
jobz = MagmaVectorsStr;
}
else if ( strcmp("-JN", argv[i]) == 0 ) {
jobz = MagmaNoVectorsStr;
}
}
if (N>0)
printf(" testing_cheevd -N %d [-JV] [-JN]\n\n", (int) N);
else {
printf("\nUsage: \n");
printf(" testing_cheevd -N %d [-JV] [-JN]\n\n", (int) N);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_cheevd -N %d [-JV] [-JN]\n\n", 1024);
N = size[7];
}
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
if ( checkres && jobz[0] == MagmaNoVectors ) {
printf( "Cannot check results when vectors are not computed (jobz='N')\n" );
checkres = false;
}
/* Query for workspace sizes */
cuFloatComplex aux_work[1];
float aux_rwork[1];
magma_int_t aux_iwork[1];
magma_cheevd( jobz[0], uplo[0],
N, h_R, N, w1,
aux_work, -1,
aux_rwork, -1,
aux_iwork, -1,
&info );
magma_int_t lwork, lrwork, liwork;
lwork = (magma_int_t) MAGMA_C_REAL( aux_work[0] );
lrwork = (magma_int_t) aux_rwork[0];
liwork = aux_iwork[0];
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, cuFloatComplex, N*N );
TESTING_MALLOC( w1, float , N );
TESTING_MALLOC( w2, float , N );
TESTING_HOSTALLOC( h_R, cuFloatComplex, N*N );
TESTING_HOSTALLOC( h_work, cuFloatComplex, lwork );
TESTING_MALLOC( rwork, float, lrwork );
TESTING_MALLOC( iwork, magma_int_t, liwork );
printf(" N CPU Time(s) GPU Time(s) \n");
printf("===================================\n");
for(i=0; i<8; i++){
if (argc==1){
N = size[i];
}
n2 = N*N;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
for( int i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_A[i*N+i], MAGMA_C_REAL(h_A[i*N+i]) );
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
/* warm up run */
magma_cheevd(jobz[0], uplo[0],
N, h_R, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
/* query for optimal workspace sizes */
magma_cheevd(jobz[0], uplo[0],
N, h_R, N, w1,
h_work, -1,
rwork, -1,
iwork, -1,
&info);
int lwork_save = lwork;
int lrwork_save = lrwork;
int liwork_save = liwork;
lwork = min( lwork, (magma_int_t) MAGMA_C_REAL( h_work[0] ));
lrwork = min( lrwork, (magma_int_t) rwork[0] );
liwork = min( liwork, iwork[0] );
//printf( "lwork %d, query %d, used %d; liwork %d, query %d, used %d\n",
// lwork_save, (magma_int_t) h_work[0], lwork,
// liwork_save, iwork[0], liwork );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_cheevd(jobz[0], uplo[0],
N, h_R, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
lwork = lwork_save;
lrwork = lrwork_save;
liwork = liwork_save;
if ( checkres ) {
/* =====================================================================
Check the results following the LAPACK's [zcds]drvst routine.
A is factored as A = U S U' and the following 3 tests computed:
(1) | A - U S U' | / ( |A| N )
(2) | I - U'U | / ( N )
(3) | S(with U) - S(w/o U) | / | S |
=================================================================== */
float temp1, temp2;
cuFloatComplex *tau;
lapackf77_chet21(&ione, uplo, &N, &izero,
h_A, &N,
w1, w1,
h_R, &N,
h_R, &N,
tau, h_work, rwork, &result[0]);
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
magma_cheevd('N', uplo[0],
N, h_R, N, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
temp1 = temp2 = 0;
for(int j=0; j<N; j++){
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / temp1;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_cheevd(jobz, uplo,
&N, h_A, &N, w2,
h_work, &lwork,
rwork, &lrwork,
iwork, &liwork,
&info);
end = get_current_time();
if (info < 0)
printf("Argument %d of cheevd had an illegal value.\n", (int) -info);
cpu_time = GetTimerValue(start,end)/1000.;
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %6.2f %6.2f\n",
(int) N, cpu_time, gpu_time);
if ( checkres ){
printf("Testing the factorization A = U S U' for correctness:\n");
printf("(1) | A - U S U' | / (|A| N) = %e\n", result[0]*eps);
printf("(2) | I - U'U | / N = %e\n", result[1]*eps);
printf("(3) | S(w/ U)-S(w/o U)|/ |S| = %e\n\n", result[2]);
}
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( h_A);
TESTING_FREE( w1);
TESTING_FREE( w2);
TESTING_FREE( rwork);
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing 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 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);
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotf2_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *d_A;
magma_int_t N, n2, lda, ldda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
gflops = FLOPS_ZPOTRF( N ) / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
magma_zsetmatrix( N, N, h_A, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zpotf2_gpu( opts.uplo, N, d_A, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotf2_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_zgetmatrix( N, N, d_A, ldda, h_R, lda );
error = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e\n",
(int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000., error );
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time*1000. );
}
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_DEVFREE( d_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssygvdx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gpu_time /*cpu_time*/;
float *h_A, *h_R, *h_B, *h_S, *h_work;
float *w1, *w2, vl=0, vu=0;
float result[2] = {0};
magma_int_t *iwork;
magma_int_t N, n2, info, il, iu, m1, m2, nb, lwork, liwork;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork;
#endif
//float d_one = 1.;
//float d_ten = 10.;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
float tolulp = opts.tolerance * lapackf77_slamch("P");
if ( opts.check && opts.jobz == MagmaNoVec ) {
fprintf( stderr, "checking results requires vectors; setting jobz=V (option -JV)\n" );
opts.jobz = MagmaVec;
}
printf(" N M GPU Time (sec)\n");
printf("============================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
n2 = N*N;
nb = magma_get_ssytrd_nb(N);
#if defined(PRECISION_z) || defined(PRECISION_c)
lwork = 2*N*nb + N*N;
lrwork = 1 + 5*N +2*N*N;
#else
lwork = 1 + 6*N*nb + 2* N*N;
#endif
liwork = 3 + 5*N;
if ( opts.fraction == 0 ) {
il = N / 10;
iu = N / 5+il;
}
else {
il = 1;
iu = (int) (opts.fraction*N);
if (iu < 1) iu = 1;
}
TESTING_MALLOC( h_A, float, n2 );
TESTING_MALLOC( h_B, float, n2 );
TESTING_MALLOC( w1, float, N );
TESTING_MALLOC( w2, float, N );
TESTING_MALLOC( iwork, magma_int_t, liwork );
TESTING_HOSTALLOC( h_R, float, n2 );
TESTING_HOSTALLOC( h_S, float, n2 );
TESTING_HOSTALLOC( h_work, float, lwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTALLOC( rwork, float, lrwork);
#endif
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
for(int i=0; i<N; i++) {
MAGMA_S_SET2REAL( h_B[i*N+i], ( MAGMA_S_REAL(h_B[i*N+i]) + 1.*N ) );
MAGMA_S_SET2REAL( h_A[i*N+i], MAGMA_S_REAL(h_A[i*N+i]) );
}
// ==================================================================
// Warmup using MAGMA
// ==================================================================
if(opts.warmup){
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_ssygvdx( opts.itype, opts.jobz, 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
gpu_time = magma_wtime();
magma_ssygvdx( opts.itype, opts.jobz, 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m1, w1,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
gpu_time = magma_wtime() - gpu_time;
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvdx routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | S(with V) - S(w/o V) | / | S |
=================================================================== */
#if defined(PRECISION_d) || defined(PRECISION_s)
float *rwork = h_work + N*N;
#endif
float temp1, temp2;
result[0] = 1.;
result[0] /= lapackf77_slansy("1", &opts.uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_slange("1", &N, &m1, h_R, &N, rwork);
if (opts.itype == 1) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_work, &N, rwork)/N;
}
else if (opts.itype == 2) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &N, rwork)/N;
}
else if (opts.itype == 3) {
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i < m1; ++i)
blasf77_sscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_ssymm("L", &opts.uplo, &N, &m1, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_slange("1", &N, &m1, h_R, &N, rwork)/N;
}
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_slacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_ssygvdx( opts.itype, 'N', 'I', opts.uplo,
N, h_R, N, h_S, N, vl, vu, il, iu, &m2, w2,
h_work, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork,
&info );
if (info != 0)
printf("magma_ssygvdx returned error %d: %s.\n",
(int) info, magma_strerror( info ));
temp1 = temp2 = 0;
for(int j=0; j < m2; j++) {
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[1] = temp2 / (((float)m2)*temp1);
}
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %5d %7.2f\n",
(int) N, (int) m1, gpu_time);
if ( opts.check ) {
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if (opts.itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
else if (opts.itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
else if (opts.itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %8.2e%s\n", result[0], (result[0] < tol ? "" : " failed"));
printf( "(2) | D(w/ Z) - D(w/o Z) | / |D| = %8.2e%s\n\n", result[1], (result[1] < tolulp ? "" : " failed"));
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( w1 );
TESTING_FREE( w2 );
#if defined(PRECISION_z) || defined(PRECISION_c)
TESTING_HOSTFREE( rwork);
#endif
TESTING_FREE( iwork );
TESTING_HOSTFREE( h_work );
TESTING_HOSTFREE( h_R );
TESTING_HOSTFREE( h_S );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgelqf
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
float flops, gpu_perf, cpu_perf;
float matnorm, work[1];
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
cuFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
/* Matrix size */
magma_int_t M = 0, N = 0, n2, lda, lwork;
magma_int_t size[10] = {1024,2048,3072,4032,5184,6016,7040,8064,9088,10112};
magma_int_t i, info, min_mn, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
N = atoi(argv[++i]);
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (N>0 && M>0)
printf(" testing_cgelqf -M %d -N %d\n\n", (int) M, (int) N);
else
{
printf("\nUsage: \n");
printf(" testing_cgelqf -M %d -N %d\n\n", (int) M, (int) N);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_cgelqf -M %d -N %d\n\n", 1024, 1024);
M = N = size[9];
}
n2 = M * N;
min_mn = min(M, N);
nb = magma_get_cgeqrf_nb(M);
TESTING_MALLOC( tau, cuFloatComplex, min_mn );
TESTING_MALLOC( h_A, cuFloatComplex, n2 );
TESTING_HOSTALLOC( h_R, cuFloatComplex, n2 );
lwork = -1;
lapackf77_cgelqf(&M, &N, h_A, &M, tau, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max( lwork, M*nb );
TESTING_HOSTALLOC( h_work, cuFloatComplex, lwork );
printf(" M N CPU GFlop/s GPU GFlop/s ||R||_F / ||A||_F\n");
printf("==========================================================\n");
for(i=0; i<10; i++){
if (argc == 1){
M = N = size[i];
}
min_mn= min(M, N);
lda = M;
n2 = lda*N;
flops = FLOPS( (float)M, (float)N ) / 1000000;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_cgelqf( M, N, h_R, lda, tau, h_work, lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of magma_cgelqf had an illegal value.\n", (int) -info);
gpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_cgelqf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
end = get_current_time();
if (info < 0)
printf("Argument %d of lapack_cgelqf had an illegal value.\n", (int) -info);
cpu_perf = flops / GetTimerValue(start, end);
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
matnorm = lapackf77_clange("f", &M, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %5d %6.2f %6.2f %e\n",
(int) M, (int) N, cpu_perf, gpu_perf,
lapackf77_clange("f", &M, &N, h_R, &lda, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
TESTING_HOSTFREE( h_work );
/* Shutdown */
TESTING_CUDA_FINALIZE();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zpotrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_R;
magma_int_t N, n2, lda, info;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double work[1], error;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("ngpu %d, uplo %c\n", (int) opts.ngpu, opts.uplo );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma - R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
n2 = lda*N;
gflops = FLOPS_ZPOTRF( N ) / 1e9;
TESTING_MALLOC( h_A, magmaDoubleComplex, n2 );
TESTING_HOSTALLOC( h_R, magmaDoubleComplex, n2 );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
magma_zmake_hpd( N, h_A, lda );
lapackf77_zlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_zpotrf( opts.uplo, N, h_R, lda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_zpotrf( &opts.uplo, &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_zpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_zlange("f", &N, &N, h_A, &lda, work);
blasf77_zaxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_zlange("f", &N, &N, h_R, &lda, work) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e%s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed") );
status |= ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE( h_A );
TESTING_HOSTFREE( h_R );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegvd
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
cuFloatComplex *h_A, *h_R, *h_B, *h_S, *h_work;
float *rwork, *w1, *w2;
magma_int_t *iwork;
float gpu_time, cpu_time;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t size[4] = {1024,2048,4100,6001};
magma_int_t i, itype, info;
magma_int_t ione = 1, izero = 0;
magma_int_t five = 5;
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.;
float d_neg_one = -1.;
float d_ten = 10.;
magma_int_t ISEED[4] = {0,0,0,1};
const char *uplo = MagmaLowerStr;
const char *jobz = MagmaVectorsStr;
itype = 1;
magma_int_t checkres;
float result[4];
int flagN = 0;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0){
N = atoi(argv[++i]);
if (N>0){
printf(" testing_chegvd -N %d\n\n", (int) N);
flagN=1;
}
else {
printf("\nUsage: \n");
printf(" testing_chegvd -N %d\n\n", (int) N);
exit(1);
}
}
if (strcmp("-itype", argv[i])==0){
itype = atoi(argv[++i]);
if (itype>0 && itype <= 3){
printf(" testing_chegvd -itype %d\n\n", (int) itype);
}
else {
printf("\nUsage: \n");
printf(" testing_chegvd -itype %d\n\n", (int) itype);
exit(1);
}
}
if (strcmp("-L", argv[i])==0){
uplo = MagmaLowerStr;
printf(" testing_chegvd -L");
}
if (strcmp("-U", argv[i])==0){
uplo = MagmaUpperStr;
printf(" testing_chegvd -U");
}
}
} else {
printf("\nUsage: \n");
printf(" testing_chegvd -L/U -N %d -itype %d\n\n", 1024, 1);
}
if(!flagN)
N = size[3];
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
n2 = N * N;
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, cuFloatComplex, n2);
TESTING_MALLOC( h_B, cuFloatComplex, n2);
TESTING_MALLOC( w1, float , N);
TESTING_MALLOC( w2, float , N);
TESTING_HOSTALLOC(h_R, cuFloatComplex, n2);
TESTING_HOSTALLOC(h_S, cuFloatComplex, n2);
magma_int_t nb = magma_get_chetrd_nb(N);
magma_int_t lwork = 2*N*nb + N*N;
magma_int_t lrwork = 1 + 5*N +2*N*N;
magma_int_t liwork = 3 + 5*N;
TESTING_HOSTALLOC(h_work, cuFloatComplex, lwork);
TESTING_MALLOC( rwork, float, lrwork);
TESTING_MALLOC( iwork, magma_int_t, liwork);
printf(" N CPU Time(s) GPU Time(s) \n");
printf("===================================\n");
for(i=0; i<4; i++){
if (!flagN){
N = size[i];
n2 = N*N;
}
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
//lapackf77_clatms( &N, &N, "U", ISEED, "P", w1, &five, &d_ten,
// &d_one, &N, &N, uplo, h_B, &N, h_work, &info);
//lapackf77_claset( "A", &N, &N, &c_zero, &c_one, h_B, &N);
lapackf77_clarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_B[i*N+i], MAGMA_C_REAL(h_B[i*N+i]) + 1.*N );
MAGMA_C_SET2REAL( h_A[i*N+i], MAGMA_C_REAL(h_A[i*N+i]) );
}
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_chegvd(itype, jobz[0], uplo[0],
N, h_R, N, h_S, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_chegvd(itype, jobz[0], uplo[0],
N, h_R, N, h_S, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvd routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | I - V V' B | / ( N ) (itype = 1,2)
| B - V V' | / ( |B| N ) (itype = 3)
(3) | S(with V) - S(w/o V) | / | S |
=================================================================== */
float temp1, temp2;
cuFloatComplex *tau;
if (itype == 1 || itype == 2){
lapackf77_claset( "A", &N, &N, &c_zero, &c_one, h_S, &N);
blasf77_cgemm("N", "C", &N, &N, &N, &c_one, h_R, &N, h_R, &N, &c_zero, h_work, &N);
blasf77_chemm("R", uplo, &N, &N, &c_neg_one, h_B, &N, h_work, &N, &c_one, h_S, &N);
result[1]= lapackf77_clange("1", &N, &N, h_S, &N, rwork) / N;
}
else if (itype == 3){
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N);
blasf77_cherk(uplo, "N", &N, &N, &d_neg_one, h_R, &N, &d_one, h_S, &N);
result[1]= lapackf77_clanhe("1",uplo, &N, h_S, &N, rwork) / N / lapackf77_clanhe("1",uplo, &N, h_B, &N, rwork);
}
result[0] = 1.;
result[0] /= lapackf77_clanhe("1",uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_clange("1",&N , &N, h_R, &N, rwork);
if (itype == 1){
blasf77_chemm("L", uplo, &N, &N, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", uplo, &N, &N, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_clange("1", &N, &N, h_work, &N, rwork)/N;
}
else if (itype == 2){
blasf77_chemm("L", uplo, &N, &N, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", uplo, &N, &N, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_clange("1", &N, &N, h_R, &N, rwork)/N;
}
else if (itype == 3){
blasf77_chemm("L", uplo, &N, &N, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", uplo, &N, &N, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_clange("1", &N, &N, h_R, &N, rwork)/N;
}
/* lapackf77_chet21(&ione, uplo, &N, &izero,
h_A, &N,
w1, w1,
h_R, &N,
h_R, &N,
tau, h_work, rwork, &result[0]);
*/
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_chegvd(itype, 'N', uplo[0],
N, h_R, N, h_S, N, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
temp1 = temp2 = 0;
for(int j=0; j<N; j++){
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / temp1;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_chegvd(&itype, jobz, uplo,
&N, h_A, &N, h_B, &N, w2,
h_work, &lwork,
rwork, &lrwork,
iwork, &liwork,
&info);
end = get_current_time();
if (info < 0)
printf("Argument %d of chegvd had an illegal value.\n", (int) -info);
cpu_time = GetTimerValue(start,end)/1000.;
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %6.2f %6.2f\n",
(int) N, cpu_time, gpu_time);
if ( checkres ){
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
if(itype==1 || itype ==2)
printf("(2) | I - Z Z' B | / N = %e\n", result[1]);
else
printf("(2) | B - Z Z' | / (|B| N) = %e\n", result[1]);
printf("(3) | D(w/ Z)-D(w/o Z)|/ |D| = %e\n\n", result[2]);
}
if (flagN)
break;
}
/* Memory clean up */
TESTING_FREE( h_A);
TESTING_FREE( h_B);
TESTING_FREE( w1);
TESTING_FREE( w2);
TESTING_FREE( rwork);
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
TESTING_HOSTFREE( h_S);
/* Shutdown */
TESTING_CUDA_FINALIZE();
}