int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time, error;
float matnorm, work[1];
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex_ptr d_lA[MagmaMaxGPUs];
/* Matrix size */
magma_int_t M = 0, N = 0, n2, n_local[4], lda, ldda, lhwork;
magma_int_t size[10] = {1000,2000,3000,4000,5000,6000,7000,8000,9000,10000};
magma_int_t i, k, nk, info, min_mn;
int max_num_gpus = 2, num_gpus = 2;
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]);
else if (strcmp("-NGPU", argv[i])==0)
num_gpus = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (M>0 && N>0)
printf(" testing_cgeqrf_gpu -M %d -N %d -NGPU %d\n\n", (int) M, (int) N, (int) num_gpus);
else
{
printf("\nUsage: \n");
printf(" testing_cgeqrf_gpu -M %d -N %d -NGPU %d\n\n",
1024, 1024, 1);
exit(1);
}
}
else {
printf("\nUsage: \n");
printf(" testing_cgeqrf_gpu -M %d -N %d -NGPU %d\n\n", 1024, 1024, 1);
M = N = size[9];
}
ldda = ((M+31)/32)*32;
n2 = M * N;
min_mn = min(M, N);
magma_int_t nb = magma_get_cgeqrf_nb(M);
if (num_gpus > max_num_gpus){
printf("More GPUs requested than available. Have to change it.\n");
num_gpus = max_num_gpus;
}
printf("Number of GPUs to be used = %d\n", (int) num_gpus);
/* Initialize */
magma_queue_t queues[MagmaMaxGPUs * 2];
magma_device_t devices[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
for(i=0;i<num_gpus;i++){
err = magma_queue_create( devices[i], &queues[2*i] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( devices[i], &queues[2*i+1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
}
/* Allocate host memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
for(i=0; i<num_gpus; i++){
n_local[i] = ((N/nb)/num_gpus)*nb;
if (i < (N/nb)%num_gpus)
n_local[i] += nb;
else if (i == (N/nb)%num_gpus)
n_local[i] += N%nb;
TESTING_MALLOC_DEV( d_lA[i], magmaFloatComplex, ldda*n_local[i] );
printf("device %2d n_local = %4d\n", (int) i, (int) n_local[i]);
}
lhwork = -1;
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, tmp, &lhwork, &info);
lhwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lhwork );
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||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;
ldda = ((M+31)/32)*32;
gflops = FLOPS( (float)M, (float)N ) * 1e-9;
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
lapackf77_cgeqrf(&M, &N, h_A, &M, tau, h_work, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
if (info < 0)
printf("Argument %d of lapack_cgeqrf had an illegal value.\n", (int) -info);
cpu_perf = gflops / cpu_time;
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
int j;
magma_queue_t *trans_queues = (magma_queue_t*)malloc(num_gpus*sizeof(magma_queue_t));
for(j=0;j<num_gpus;j++){
trans_queues[j] = queues[2*j];
}
// warm-up
magmablas_csetmatrix_1D_bcyclic(M, N, h_R, lda, d_lA, ldda, num_gpus, nb, trans_queues);
magma_cgeqrf2_mgpu( num_gpus, M, N, d_lA, ldda, tau, &info, queues);
magmablas_csetmatrix_1D_bcyclic(M, N, h_R, lda, d_lA, ldda, num_gpus, nb, trans_queues);
gpu_time = magma_wtime();
magma_cgeqrf2_mgpu( num_gpus, M, N, d_lA, ldda, tau, &info, queues);
gpu_time = magma_wtime() - gpu_time;
if (info < 0)
printf("Argument %d of magma_cgeqrf2 had an illegal value.\n", (int) -info);
gpu_perf = gflops / gpu_time;
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magmablas_cgetmatrix_1D_bcyclic(M, N, d_lA, ldda, h_R, lda, num_gpus, nb, trans_queues);
matnorm = lapackf77_clange("f", &M, &N, h_A, &M, work);
blasf77_caxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
printf("%5d %5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
lapackf77_clange("f", &M, &N, h_R, &M, work) / matnorm);
if (argc != 1)
break;
}
/* Memory clean up */
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_A );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_R );
for(i=0; i<num_gpus; i++){
TESTING_FREE_DEV( d_lA[i] );
magma_queue_destroy(queues[2*i]);
magma_queue_destroy(queues[2*i+1]);
}
/* Shutdown */
magma_finalize();
}
int main(int argc, char **argv)
{
#if (GPUSHMEM >= 200)
TESTING_CUDA_INIT();
cudaSetDevice(0);
magma_timestr_t start, end;
float flops, magma_perf, cuda_perf, error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t n_local[4];
FILE *fp ;
magma_int_t N, m, i, j, lda, LDA, M;
magma_int_t matsize;
magma_int_t vecsize;
magma_int_t istart = 64;
magma_int_t incx = 1;
char uplo = MagmaLower;
cuFloatComplex alpha = MAGMA_C_MAKE(1., 0.); // MAGMA_C_MAKE( 1.5, -2.3 );
cuFloatComplex beta = MAGMA_C_MAKE(0., 0.); // MAGMA_C_MAKE( -0.6, 0.8 );
cuFloatComplex *A, *X, *Y[4], *Ycublas, *Ymagma;
cuFloatComplex *dA, *dX[4], *dY[4], *d_lA[4], *dYcublas ;
cudaStream_t stream[4][10];
cuFloatComplex *C_work;
cuFloatComplex *dC_work[4];
magma_int_t num_gpus = 1, max_num_gpus, nb;
magma_int_t blocks, workspace;
magma_int_t offset;
// offset = 257;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0)
{
N = atoi(argv[++i]);
istart = N;
}
else if (strcmp("-M", argv[i])==0)
M = atoi(argv[++i]);
else if (strcmp("-NGPU", argv[i])==0)
num_gpus = atoi(argv[++i]);
else if (strcmp("-offset", argv[i])==0)
offset = atoi(argv[++i]);
}
if ( M == 0 ) {
M = N;
}
if ( N == 0 ) {
N = M;
}
if (M>0 && N>0)
{ printf(" testing_chemv_mgpu -M %d -N %d -NGPU %d\n\n", M, N, num_gpus);
printf(" in %c side \n", uplo);
}
else
{
printf("\nUsage: \n");
printf(" testing_chemv_mgpu -M %d -N %d -NGPU %d\n\n",
1024, 1024, 1);
exit(1);
}
}
else {
#if defined(PRECISION_z)
M = N = 8000;
#else
M = N = 12480;
#endif
num_gpus = 2;
offset = 0;
printf("\nUsage: \n");
printf(" testing_chemv_mgpu -M %d -N %d -NGPU %d\n\n", M, N, num_gpus);
}
///////////////////////////////////////////////////////////////////////////////////////
cudaGetDeviceCount(&max_num_gpus);
if (num_gpus > max_num_gpus){
printf("More GPUs requested than available. Have to change it.\n");
num_gpus = max_num_gpus;
}
printf("Number of GPUs to be used = %d\n", num_gpus);
for(int i=0; i< num_gpus; i++)
{
cudaStreamCreate(&stream[i][0]);
}
LDA = ((N+31)/32)*32;
matsize = N*LDA;
vecsize = N*incx;
nb = 32;
//nb = 64;
printf("block size = %d\n", nb);
TESTING_MALLOC( A, cuFloatComplex, matsize );
TESTING_MALLOC( X, cuFloatComplex, vecsize );
TESTING_MALLOC( Ycublas, cuFloatComplex, vecsize );
TESTING_MALLOC( Ymagma, cuFloatComplex, vecsize );
for(i=0; i<num_gpus; i++)
{
TESTING_MALLOC( Y[i], cuFloatComplex, vecsize );
}
cudaSetDevice(0);
TESTING_DEVALLOC( dA, cuFloatComplex, matsize );
TESTING_DEVALLOC( dYcublas, cuFloatComplex, vecsize );
for(i=0; i<num_gpus; i++)
{
n_local[i] = ((N/nb)/num_gpus)*nb;
if (i < (N/nb)%num_gpus)
n_local[i] += nb;
else if (i == (N/nb)%num_gpus)
n_local[i] += N%nb;
cudaSetDevice(i);
TESTING_DEVALLOC( d_lA[i], cuFloatComplex, LDA*n_local[i] );// potentially bugged
TESTING_DEVALLOC( dX[i], cuFloatComplex, vecsize );
TESTING_DEVALLOC( dY[i], cuFloatComplex, vecsize );
printf("device %2d n_local = %4d\n", i, n_local[i]);
}
cudaSetDevice(0);
///////////////////////////////////////////////////////////////////////
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &matsize, A );
/* Make A hermitian */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
A[i*LDA+i] = MAGMA_C_MAKE( MAGMA_C_REAL(A[i*LDA+i]), 0. );
for(j=0; j<i; j++)
A[i*LDA+j] = cuConjf(A[j*LDA+i]);
}
}
blocks = N / nb + (N % nb != 0);
workspace = LDA * (blocks + 1);
TESTING_MALLOC( C_work, cuFloatComplex, workspace );
for(i=0; i<num_gpus; i++){
cudaSetDevice(i);
TESTING_DEVALLOC( dC_work[i], cuFloatComplex, workspace );
//fillZero(dC_work[i], workspace);
}
cudaSetDevice(0);
//////////////////////////////////////////////////////////////////////////////////////////////
fp = fopen ("results_chemv_mgpu.csv", "w") ;
if( fp == NULL ){ printf("Couldn't open output file\n"); exit(1);}
printf("HEMV cuFloatComplex Precision\n\n");
printf( " n CUBLAS,Gflop/s MAGMABLAS,Gflop/s \"error\"\n"
"==============================================================\n");
fprintf(fp, " n CUBLAS,Gflop/s MAGMABLAS,Gflop/s \"error\"\n"
"==============================================================\n");
// for( offset = 0; offset< N; offset ++ )
for(int size = istart ; size <= N ; size += 128)
{
// printf("offset = %d ", offset);
m = size ;
// m = N;
// lda = ((m+31)/32)*32;//
lda = LDA;
flops = FLOPS( (float)m ) / 1e6;
printf( "N %5d ", m );
fprintf( fp, "%5d, ", m );
vecsize = m * incx;
lapackf77_clarnv( &ione, ISEED, &vecsize, X );
lapackf77_clarnv( &ione, ISEED, &vecsize, Y[0] );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
cudaSetDevice(0);
magmablas_csetmatrix_1D_bcyclic(m, m, A, LDA, d_lA, lda, num_gpus, nb);
cudaSetDevice(0);
magma_csetmatrix( m, m, A, LDA, dA, lda );
magma_csetvector( m, Y[0], incx, dYcublas, incx );
for(i=0; i<num_gpus; i++){
cudaSetDevice(i);
magma_csetvector( m, X, incx, dX[i], incx );
magma_csetvector( m, Y[0], incx, dY[i], incx );
blocks = m / nb + (m % nb != 0);
magma_csetmatrix( lda, blocks, C_work, LDA, dC_work[i], lda );
//cudaMemset(dC_work[i], 0, sizeof( cuFloatComplex) * lda * blocks);
}
cudaSetDevice(0);
start = get_current_time();
cublasChemv( uplo, m-offset, alpha, dA + offset + offset * lda, lda, dX[0] + offset, incx, beta, dYcublas + offset, incx );
end = get_current_time();
magma_cgetvector( m, dYcublas, incx, Ycublas, incx );
cuda_perf = flops / GetTimerValue(start,end);
printf( "%11.2f", cuda_perf );
fprintf(fp, "%11.2f,", cuda_perf );
cudaSetDevice(0);
start = get_current_time();
if(nb == 32)
{
magmablas_chemv2_mgpu_32_offset( uplo, m, alpha, d_lA, lda, dX, incx, beta, dY, incx,
dC_work, workspace, num_gpus, nb, offset);
}
else // nb = 64
{
magmablas_chemv2_mgpu_offset( uplo, m, alpha, d_lA, lda, dX, incx, beta, dY, incx,
dC_work, workspace, num_gpus, nb, offset);
}
for(i=1; i<num_gpus; i++)
{
cudaSetDevice(i);
cudaDeviceSynchronize();
}
end = get_current_time();
magma_perf = flops / GetTimerValue(start,end);
printf( "%11.2f", magma_perf );
fprintf(fp, "%11.2f,", magma_perf );
for(i=0; i<num_gpus; i++)
{
cudaSetDevice(i);
magma_cgetvector( m, dY[i], incx, Y[i], incx );
}
cudaSetDevice(0);
#ifdef validate
for( j= offset;j<m;j++)
{
for(i=1; i<num_gpus; i++)
{
// printf("Y[%d][%d] = %15.14f\n", i, j, Y[i][j].x);
#if defined(PRECISION_z) || defined(PRECISION_c)
Y[0][j].x = Y[0][j].x + Y[i][j].x;
Y[0][j].y = Y[0][j].y + Y[i][j].y;
#else
Y[0][j] = Y[0][j] + Y[i][j];
#endif
}
}
/*
#if defined(PRECISION_z) || defined(PRECISION_c)
for( j=offset;j<m;j++)
{
if(Y[0][j].x != Ycublas[j].x)
{
printf("Y-multi[%d] = %f, %f\n", j, Y[0][j].x, Y[0][j].y );
printf("Ycublas[%d] = %f, %f\n", j, Ycublas[j].x, Ycublas[j].y);
}
}
#else
for( j=offset;j<m;j++)
{
if(Y[0][j] != Ycublas[j])
{
printf("Y-multi[%d] = %f\n", j, Y[0][j] );
printf("Ycublas[%d] = %f\n", j, Ycublas[j]);
}
}
#endif
*/
/* =====================================================================
Computing the Difference Cublas VS Magma
=================================================================== */
int nw = m - offset ;
blasf77_caxpy( &nw, &c_neg_one, Y[0] + offset, &incx, Ycublas + offset, &incx);
error = lapackf77_clange( "M", &nw, &ione, Ycublas + offset, &nw, work );
#if 0
printf( "\t\t %8.6e", error / m );
fprintf( fp, "\t\t %8.6e", error / m );
/*
* Extra check with cblas vs magma
*/
cblas_ccopy( m, Y, incx, Ycublas, incx );
cblas_chemv( CblasColMajor, CblasLower, m,
CBLAS_SADDR(alpha), A, LDA, X, incx,
CBLAS_SADDR(beta), Ycublas, incx );
blasf77_caxpy( &m, &c_neg_one, Ymagma, &incx, Ycublas, &incx);
error = lapackf77_clange( "M", &m, &ione, Ycublas, &m, work );
#endif
printf( "\t\t %8.6e", error / m );
fprintf( fp, "\t\t %8.6e", error / m );
#endif
printf("\n");
fprintf(fp, "\n");
}
fclose( fp ) ;
/* Free Memory */
TESTING_FREE( A );
TESTING_FREE( X );
TESTING_FREE( Ycublas );
TESTING_FREE( Ymagma );
TESTING_FREE( C_work );
TESTING_DEVFREE( dA );
TESTING_DEVFREE( dYcublas );
for(i=0; i<num_gpus; i++)
{
TESTING_FREE( Y[i] );
cudaSetDevice(i);
TESTING_DEVFREE( d_lA[i] )
TESTING_DEVFREE( dX[i] );
TESTING_DEVFREE( dY[i] );
TESTING_DEVFREE( dC_work[i] );
}
cudaSetDevice(0);
///////////////////////////////////////////////////////////
/* Free device */
TESTING_CUDA_FINALIZE();
#endif
return 0;
}