/* ////////////////////////////////////////////////////////////////////////////
-- 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;
}
/**
Purpose
-------
DGEHRD reduces a DOUBLE_PRECISION general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H . This version
stores the triangular matrices used in the factorization so that they can
be applied directly (i.e., without being recomputed) later. As a result,
the application of Q is much faster.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ilo INTEGER
@param[in]
ihi INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to DGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, and the
elements below the first subdiagonal, with the array TAU,
represent the orthogonal matrix Q as a product of elementary
reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE_PRECISION array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= max(1,N).
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
T DOUBLE_PRECISION array, dimension NB*N,
where NB is the optimal blocksize. It stores the NB*NB blocks
of the triangular T matrices used in the reduction.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
---------------
The matrix Q is represented as a product of (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
@verbatim
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
@endverbatim
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
This version stores the T matrices, for later use in magma_dorghr.
@ingroup magma_dgeev_comp
********************************************************************/
extern "C" magma_int_t
magma_dgehrd_m(
magma_int_t n, magma_int_t ilo, magma_int_t ihi,
double *A, magma_int_t lda,
double *tau,
double *work, magma_int_t lwork,
double *T,
magma_int_t *info)
{
#define A( i, j ) (A + (i) + (j)*lda)
#define dA( d, i, j ) (data.A[d] + (i) + (j)*ldda)
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
magma_int_t nb = magma_get_dgehrd_nb(n);
magma_int_t nh, iws, ldda, min_lblocks, max_lblocks, last_dev, d;
magma_int_t dpanel, di, nlocal, i, i2, ib, ldwork;
magma_int_t iinfo;
magma_int_t lquery;
struct dgehrd_data data;
int ngpu = magma_num_gpus();
*info = 0;
iws = n*(nb + nb*ngpu);
work[0] = MAGMA_D_MAKE( iws, 0 );
lquery = (lwork == -1);
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
// Adjust from 1-based indexing
ilo -= 1;
// Quick return if possible
nh = ihi - ilo;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// Set elements 0:ILO-1 and IHI-1:N-2 of TAU to zero
for (i = 0; i < ilo; ++i)
tau[i] = c_zero;
for (i = max(0,ihi-1); i < n-1; ++i)
tau[i] = c_zero;
// set T to zero
lapackf77_dlaset( "Full", &nb, &n, &c_zero, &c_zero, T, &nb );
// set to null, to simplify cleanup code
for( d = 0; d < ngpu; ++d ) {
data.A[d] = NULL;
data.streams[d] = NULL;
}
// If not enough workspace, use unblocked code
if ( lwork < iws ) {
nb = 1;
}
if (nb == 1 || nb >= nh) {
// Use unblocked code below
i = ilo;
}
else {
// Use blocked code
// allocate memory on GPUs for A and workspaces
ldda = ((n+31)/32)*32;
min_lblocks = (n / nb) / ngpu;
max_lblocks = ((n-1) / nb) / ngpu + 1;
last_dev = (n / nb) % ngpu;
// V and Vd need to be padded for copying in mdlahr2
data.ngpu = ngpu;
data.ldda = ldda;
data.ldv = nb*max_lblocks*ngpu;
data.ldvd = nb*max_lblocks;
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
nlocal = min_lblocks*nb;
if ( d < last_dev ) {
nlocal += nb;
}
else if ( d == last_dev ) {
nlocal += (n % nb);
}
ldwork = nlocal*ldda // A
+ nb*data.ldv // V
+ nb*data.ldvd // Vd
+ nb*ldda // Y
+ nb*ldda // W
+ nb*nb; // Ti
if ( MAGMA_SUCCESS != magma_dmalloc( &data.A[d], ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto CLEANUP;
}
data.V [d] = data.A [d] + nlocal*ldda;
data.Vd[d] = data.V [d] + nb*data.ldv;
data.Y [d] = data.Vd[d] + nb*data.ldvd;
data.W [d] = data.Y [d] + nb*ldda;
data.Ti[d] = data.W [d] + nb*ldda;
magma_queue_create( &data.streams[d] );
}
// Copy the matrix to GPUs
magma_dsetmatrix_1D_col_bcyclic( n, n, A, lda, data.A, ldda, ngpu, nb );
// round ilo down to block boundary
ilo = (ilo/nb)*nb;
for (i = ilo; i < ihi - 1 - nb; i += nb) {
// Reduce columns i:i+nb-1 to Hessenberg form, returning the
// matrices V and T of the block reflector H = I - V*T*V'
// which performs the reduction, and also the matrix Y = A*V*T
// Get the current panel (no need for the 1st iteration)
dpanel = (i / nb) % ngpu;
di = ((i / nb) / ngpu) * nb;
if ( i > ilo ) {
magma_setdevice( dpanel );
magma_dgetmatrix( ihi-i, nb,
dA(dpanel, i, di), ldda,
A(i,i), lda );
}
// add 1 to i for 1-based index
magma_dlahr2_m( ihi, i+1, nb, A(0,i), lda,
&tau[i], &T[i*nb], nb, work, n, &data );
magma_dlahru_m( n, ihi, i, nb, A, lda, &data );
// copy first i rows above panel to host
magma_setdevice( dpanel );
magma_dgetmatrix_async( i, nb,
dA(dpanel, 0, di), ldda,
A(0,i), lda, data.streams[dpanel] );
}
// Copy remainder to host, block-by-block
for( i2 = i; i2 < n; i2 += nb ) {
ib = min( nb, n-i2 );
d = (i2 / nb) % ngpu;
di = (i2 / nb) / ngpu * nb;
magma_setdevice( d );
magma_dgetmatrix( n, ib,
dA(d, 0, di), ldda,
A(0,i2), lda );
}
}
// Use unblocked code to reduce the rest of the matrix
// add 1 to i for 1-based index
i += 1;
lapackf77_dgehd2(&n, &i, &ihi, A, &lda, tau, work, &iinfo);
work[0] = MAGMA_D_MAKE( iws, 0 );
CLEANUP:
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_free( data.A[d] );
magma_queue_destroy( data.streams[d] );
}
magma_setdevice( orig_dev );
return *info;
} /* magma_dgehrd */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgetrf_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;
double *h_A;
double *d_lA[ MagmaMaxGPUs ];
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, n_local, ngpu;
magma_int_t info, min_mn, nb, ldn_local;
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 itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_dgetrf_nb( M );
gflops = FLOPS_DGETRF( 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 );
}
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
// Allocate device memory
for( int dev=0; dev < ngpu; dev++){
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
ldn_local = ((n_local+31)/32)*32; // TODO why?
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], double, ldda*ldn_local );
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_dgetrf( &M, &N, h_A, &lda, ipiv, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( M, N, h_A, lda );
magma_dsetmatrix_1D_col_bcyclic( M, N, h_A, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_dgetrf_mgpu( ngpu, M, N, d_lA, ldda, ipiv, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgetrf_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_A, lda, ngpu, nb );
/* =====================================================================
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 ? "ok" : "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 ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- 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];
magmaDouble_ptr 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;
opts.parse_opts( argc, argv );
opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code
opts.lapack |= (opts.check == 2); // check (-c2) implies lapack (-l)
magma_int_t status = 0;
double eps = lapackf77_dlamch("E");
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("%%===============================================================================================\n");
}
else {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) ||R||_F /(M*||A||_F)\n");
printf("%%=========================================================================\n");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
nb = magma_get_dgeqrf_nb( M, N );
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
// ngpu must be at least the number of blocks
ngpu = min( opts.ngpu, magma_ceildiv(N,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, NULL, &M, NULL, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_D_REAL( tmp[0] );
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_work, double, lhwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
// Allocate device memory
for( int dev = 0; dev < ngpu; dev++ ) {
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], double, ldda*n_local );
}
/* Initialize the matrix */
for( int j=0; j < 4; j++ )
ISEED2[j] = ISEED[j]; // save seeds
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
double *tau2;
TESTING_MALLOC_CPU( tau2, double, min_mn );
cpu_time = magma_wtime();
lapackf77_dgeqrf( &M, &N, h_A, &lda, tau2, 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_CPU( tau2 );
}
/* ====================================================================
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 );
if ( opts.check == 1 && M >= N ) {
/* =====================================================================
Check the result -- dqrt02 requires M >= N
=================================================================== */
magma_int_t lwork = n2+N;
double *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC_CPU( h_W1, double, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, double, n2 ); // R
TESTING_MALLOC_CPU( h_W3, double, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, double, M ); // RWORK
lapackf77_dlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_dqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0], results[1] );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0], results[1] );
}
// todo also check results[1] < tol?
printf(" %s\n", (results[0] < tol ? "ok" : "failed"));
status += ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
}
else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work );
blasf77_daxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work ) / (min_mn*error);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) ---",
(int) M, (int) N, gpu_perf, gpu_time);
}
printf("%s\n", (opts.check != 0 ? " (error check only for M >= N)" : ""));
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgetrf_mgpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
real_Double_t gpu_perf1, gpu_time1, gpu_perf2, gpu_time2, gpu_perf3, gpu_time3, alloc_time, free_time;
double error;
double *h_A;
double *d_lA[ MagmaMaxGPUs ];
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, n_local, ngpu, NB;
magma_int_t info, min_mn, nb, ldn_local;
magma_int_t status = 0;
magma_int_t P=-1; /*Number of threads in the CPU part*/
double d_cpu=-1; /*pourcentgae of the matrix to allocate in the cpu part*/
magma_int_t Pr=-1; /*Number of threads for the panel*/
magma_int_t async_nb; /*Block size*/
double *WORK;
magma_int_t WORK_LD, WORK_n;
double **dlpanelT;
magma_int_t dlpanelT_m, dlpanelT_n;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
P = opts.nthread;
async_nb = opts.nb;
Pr = opts.panel_nthread;
d_cpu = 0.0;
#if defined(CPU_PEAK) && defined(GPU_PEAK)
d_cpu = magma_amc_recommanded_dcpu(opts.nthread, CPU_PEAK, opts.ngpu, GPU_PEAK);
#endif
if(opts.fraction_dcpu!=0){ /*Overwrite the one computed with the model*/
d_cpu = opts.fraction_dcpu;
}
magma_assert(d_cpu > 0 && d_cpu<=1.0,
"error: The cpu fraction is invalid. Ensure you use --fraction_dcpu with fraction_dcpu in [0.0, 1.0] or compile with both -DCPU_PEAK=<cpu peak performance> and -DGPU_PEAK=<gpu peak performance> set.\n");
printf("Asynchronous recursif LU... nb:%d, nbcores:%d, dcpu:%f, panel_nbcores:%d, ngpu: %d\n", async_nb, P, d_cpu, Pr, opts.ngpu);
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) GPU_Async_v2 GFlop/s (sec) GPU_Async_work_v2 GFlop/s (sec)");
if ( opts.check == 2 ) {
printf(" |Ax-b|/(N*|A|*|x|)\n");
}
else {
printf(" |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;
//nb = magma_get_dgetrf_nb( M );
gflops = FLOPS_DGETRF( M, N ) / 1e9;
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
/*set default number of threads for lapack*/
magma_setlapack_numthreads(P);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_dgetrf( &M, &N, h_A, &lda, ipiv, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
nb = magma_get_dgetrf_nb( M );
// 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 );
}
// 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;
ldn_local = ((n_local+31)/32)*32; // TODO why?
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], double, ldda*ldn_local );
}
init_matrix( M, N, h_A, lda );
magma_dsetmatrix_1D_col_bcyclic( M, N, h_A, lda, d_lA, ldda, ngpu, nb );
gpu_time1 = magma_wtime();
magma_dgetrf_mgpu( ngpu, M, N, d_lA, ldda, ipiv, &info );
gpu_time1 = magma_wtime() - gpu_time1;
gpu_perf1 = gflops / gpu_time1;
if (info != 0)
printf("magma_dgetrf_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_A, lda, ngpu, nb );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
/* ====================================================================
Performs operation using MAGMA_Async: This interface allocate workspace internally
=================================================================== */
/*For the benchmark we have 2 approaches*/
/*1. use directly magma_amc */
/*2. use magma_amc_work and add pinned memory time*/
/*We choose approach 2*/
/*
nb = async_nb;
// 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 );
}
// Allocate device memory
n_local = numcols2p(0, N, nb, ngpu);
ldn_local = n_local;
//ldn_local = ((n_local+31)/32)*32;
for( int dev=0; dev < ngpu; dev++){
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], double, ldda*ldn_local );
}
init_matrix( M, N, h_A, lda );
magma_dsetmatrix_1D_col_bcyclic( M, N, h_A, lda, d_lA, ldda, ngpu, nb );
// Switch to the sequential version of BLAS
magma_setlapack_numthreads(1);
magma_amc_init(P, d_cpu, Pr, nb);
gpu_time2 = magma_wtime();
magma_dgetrf_async_mgpu( ngpu, M, N, d_lA, ldda, ipiv, &info );
gpu_time2 = magma_wtime() - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
magma_amc_finalize();
if (info != 0)
printf("magma_dgetrf_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_A, lda, ngpu, nb );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
*/
/* ====================================================================
Performs operation using MAGMA_Async_Work
=================================================================== */
nb = async_nb;
// 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 );
}
// Allocate device memory
n_local = numcols2p(0, N, nb, ngpu);
ldn_local = n_local;
//ldn_local = ((n_local+31)/32)*32;
for( int dev=0; dev < ngpu; dev++){
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], double, ldda*ldn_local );
}
init_matrix( M, N, h_A, lda );
magma_dsetmatrix_1D_col_bcyclic( M, N, h_A, lda, d_lA, ldda, ngpu, nb );
// Switch to the sequential version of BLAS
magma_setlapack_numthreads(1);
//Compute workspace dimension
WORK_LD = M;
NB = (int) ceil( (double) N / nb);
WORK_n = (int) ceil(N*d_cpu)+nb; /*TODO:remove +nb replace with A_N*/
//WORK_n = NSplit(NB, d_cpu)*nb;
if(WORK_n<nb) WORK_n = nb;//make sure workspace has at least one block column
//Make LD and n multiple of 32
//if(WORK_LD%32!=0) WORK_LD = ((WORK_LD + 31)/32)*32;
//if(WORK_n%32!=0) WORK_n = ((WORK_n + 31)/32)*32;
//Allocate workspace
alloc_time = magma_wtime();
if (MAGMA_SUCCESS != magma_dmalloc_pinned(&WORK, WORK_LD*WORK_n)) {
//if (MAGMA_SUCCESS != magma_dmalloc_cpu(&WORK, WORK_LD*WORK_n)) {
info = MAGMA_ERR_HOST_ALLOC;
printf("magma_dmalloc_pinned returned error %d: %s.\n ", (int) info);
}
/* Workspace for the panels on the GPU*/
dlpanelT_m = WORK_n; /*assume that the cpu and gpu use the same buffer size*/
dlpanelT_n = M;
dlpanelT = (double **) malloc(ngpu*sizeof(double*));
for(int dev=0;dev<ngpu;dev++){
magma_setdevice(dev);
if (MAGMA_SUCCESS != magma_dmalloc(&dlpanelT[dev], dlpanelT_m*dlpanelT_n)) {
info = MAGMA_ERR_DEVICE_ALLOC;
printf("magma_dmalloc returned error %d: %s.\n ", (int) info);
}
}
alloc_time = magma_wtime() - alloc_time;
//First touch the workspace with each thread. This may be needed to avoid using numactl --interleave
//magma_amc_dmemset(WORK, 0.0, WORK_LD*WORK_n, 256, P); //nb
//#pragma omp parallel for private(info) schedule(static,nb)
//for(info=0;info<WORK_LD*WORK_n;info++) WORK[info] = 0.0; //alternative first touch by the thread
magma_amc_init(P, d_cpu, Pr, nb);
gpu_time3 = magma_wtime();
magma_dgetrf_mgpu_work_amc_v3(ngpu, M, N, d_lA, ldda, ipiv, &info, WORK, WORK_LD, WORK_n);
gpu_time3 = magma_wtime() - gpu_time3;
gpu_perf3 = gflops / gpu_time3;
magma_amc_finalize();
if (info != 0)
printf("magma_dgetrf_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_A, lda, ngpu, nb );
//Free workspace
free_time = magma_wtime();
magma_free_pinned(WORK);
for(int dev=0;dev<ngpu;dev++){
magma_setdevice(dev);
magma_free(dlpanelT[dev]);
}
free(dlpanelT);
free_time = magma_wtime() - free_time;
/*DEDUCE t2, JUST FOR THE BENCHMARK*/
gpu_time2 = gpu_time3 + alloc_time + free_time;
gpu_perf2 = gflops / gpu_time2;
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
/* =====================================================================
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 );
}
*/
printf("%5d %5d", (int) M, (int) N);
if(cpu_perf!=0.0){
printf(" %7.2f (%7.2f)", cpu_perf, cpu_time);
}
else{
printf(" --- ( --- )");
}
if(gpu_perf1!=0.0){
printf(" %7.2f (%7.2f)", gpu_perf1, gpu_time1);
}
else{
printf(" --- ( --- )");
}
if(gpu_perf2!=0.0){
printf(" %7.2f (%7.2f)", gpu_perf2, gpu_time2);
}
else{
printf(" --- ( --- )");
}
if(gpu_perf3!=0.0){
printf(" %7.2f (%7.2f)", gpu_perf3, gpu_time3);
}
else{
printf(" --- ( --- )");
}
if ( opts.check == 2 ) {
error = get_residual( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else if ( opts.check ) {
error = get_LU_error( M, N, h_A, lda, ipiv );
printf(" %8.2e%s\n", error, (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dpotrf_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;
double *d_lA[ MagmaMaxGPUs ];
magma_int_t N, n2, lda, ldda, max_size, ngpu;
magma_int_t info, nb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("ngpu = %d, uplo = %s\n", (int) opts.ngpu, lapack_uplo_const(opts.uplo) );
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
nb = magma_get_dpotrf_nb( N );
gflops = FLOPS_DPOTRF( 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 );
}
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
// Allocate device memory
// matrix is distributed by block-rows or block-columns
// this is maximum size that any GPU stores;
// size is rounded up to full blocks in both rows and columns
max_size = nb*(1+N/(nb*ngpu)) * nb*((N+nb-1)/nb);
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], double, max_size );
}
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
magma_dmake_hpd( N, h_A, lda );
lapackf77_dlacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dpotrf( lapack_uplo_const(opts.uplo), &N, h_A, &lda, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
if ( opts.uplo == MagmaUpper ) {
ldda = ((N+nb-1)/nb)*nb;
magma_dsetmatrix_1D_col_bcyclic( N, N, h_R, lda, d_lA, ldda, ngpu, nb );
} else {
ldda = (1+N/(nb*ngpu))*nb;
magma_dsetmatrix_1D_row_bcyclic( N, N, h_R, lda, d_lA, ldda, ngpu, nb );
}
gpu_time = magma_wtime();
magma_dpotrf_mgpu( ngpu, opts.uplo, N, d_lA, ldda, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dpotrf_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.uplo == MagmaUpper ) {
magma_dgetmatrix_1D_col_bcyclic( N, N, d_lA, ldda, h_R, lda, ngpu, nb );
} else {
magma_dgetmatrix_1D_row_bcyclic( N, N, d_lA, ldda, h_R, lda, ngpu, nb );
}
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
for( int dev=0; dev < ngpu; dev++ ){
magma_setdevice( dev );
magma_device_sync();
}
if ( opts.lapack ) {
error = lapackf77_dlange("f", &N, &N, h_A, &lda, work );
blasf77_daxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_dlange("f", &N, &N, h_R, &lda, work ) / error;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
}
else {
printf("%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
for( int dev=0; dev < ngpu; dev++ ){
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
DGEQRF4 computes a QR factorization of a DOUBLE_PRECISION M-by-N matrix A:
A = Q * R using multiple GPUs. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
Arguments
---------
@param[in]
num_gpus INTEGER
The number of GPUs to be used for the factorization.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau DOUBLE_PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_dgeqrf_nb(M).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_dgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_dgeqrf4(magma_int_t num_gpus, magma_int_t m, magma_int_t n,
double *A, magma_int_t lda, double *tau,
double *work, magma_int_t lwork,
magma_int_t *info )
{
double *da[MagmaMaxGPUs];
double c_one = MAGMA_D_ONE;
int i, k, ldda;
*info = 0;
int nb = magma_get_dgeqrf_nb(min(m, n));
int lwkopt = n * nb;
work[0] = MAGMA_D_MAKE( (double)lwkopt, 0 );
int lquery = (lwork == -1);
if (num_gpus < 0 || num_gpus > MagmaMaxGPUs) {
*info = -1;
} else if (m < 0) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
ldda = ((m+31)/32)*32;
magma_int_t n_local[MagmaMaxGPUs];
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;
magma_setdevice(i);
// TODO on failure, free previously allocated memory
if (MAGMA_SUCCESS != magma_dmalloc( &da[i], ldda*n_local[i] )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (m > nb && n > nb) {
/* Copy the matrix to the GPUs in 1D block cyclic distribution */
magma_dsetmatrix_1D_col_bcyclic(m, n, A, lda, da, ldda, num_gpus, nb);
/* Factor using the GPU interface */
magma_dgeqrf2_mgpu( num_gpus, m, n, da, ldda, tau, info);
/* Copy the matrix back from the GPUs to the CPU */
magma_dgetmatrix_1D_col_bcyclic(m, n, da, ldda, A, lda, num_gpus, nb);
}
else {
lapackf77_dgeqrf(&m, &n, A, &lda, tau, work, &lwork, info);
}
/* Free the allocated GPU memory */
for (i=0; i < num_gpus; i++) {
magma_setdevice(i);
magma_free( da[i] );
}
magma_setdevice( orig_dev );
return *info;
} /* magma_dgeqrf4 */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_dsymm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 3.456, 5.678 );
double beta = MAGMA_D_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double Anorm, error, work[1];
double *hA, *hB, *hC, *hR;
magmaDouble_ptr dA[MagmaMaxGPUs], dB[MagmaMaxGPUs], dC[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magmaDouble_ptr dA2;
magma_int_t i, j, dev, M, N, size, lda, ldb, ldc, ldda, lddb, lddc, msize, nb;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t ncmplx = 0;
magma_buildconnection_mgpu( gnode, &ncmplx, opts.ngpu );
printf("%% Initializing communication pattern... GPU-ncmplx %d\n", (int) ncmplx);
for (i=0; i < ncmplx; ++i) {
magma_int_t myngpu = gnode[i][MagmaMaxGPUs];
printf("%% cmplx %d has %d GPUs:", i, myngpu);
for (j=0; j < myngpu; ++j) {
printf(" %d", (int) gnode[i][j]);
if (j < myngpu-1) {
printf(",");
}
}
printf("\n");
}
// number of queues per GPU. Requires ngpu.
magma_int_t nqueue = opts.ngpu;
// number of events per GPU. Require ngpu*ngpu.
magma_int_t nevents = opts.ngpu*opts.ngpu;
magma_queue_t queues[MagmaMaxGPUs][20], queues0[MagmaMaxGPUs];
magma_event_t events[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs + 10];
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( i = 0; i < nqueue; ++i ) {
magma_queue_create( dev, &queues[dev][i] );
}
queues0[dev] = queues[dev][0];
for( i = 0; i < nevents; ++i ) {
cudaEventCreateWithFlags( &events[dev][i], cudaEventDisableTiming );
}
}
printf("%% nb %d, ngpu %d, version %d\n", (int) nb, (int) opts.ngpu, (int) opts.version );
printf("%% M N nb offset CPU Gflop/s (sec) GPU Gflop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||B||\n");
printf("%%========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M; // TODO depends on side
ldb = M;
ldc = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
lddc = magma_roundup( ldc, opts.align ); // multiple of 32 by default
gflops = FLOPS_DSYMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = lddc*N + (M*N)*opts.ngpu;
TESTING_MALLOC_CPU( hA, double, lda*M );
TESTING_MALLOC_CPU( hB, double, ldb*N );
TESTING_MALLOC_CPU( hC, double, ldc*N );
TESTING_MALLOC_PIN( hR, double, ldc*N );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( dA[dev], double, ldda*mlocal );
TESTING_MALLOC_DEV( dB[dev], double, lddb*N );
TESTING_MALLOC_DEV( dC[dev], double, lddc*N );
TESTING_MALLOC_DEV( dwork[dev], double, dworksiz );
}
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, double, ldda*M );
}
size = lda*M;
lapackf77_dlarnv( &ione, iseed, &size, hA );
magma_dmake_symmetric( M, hA, lda );
size = ldb*N;
lapackf77_dlarnv( &ione, iseed, &size, hB );
size = ldc*N;
lapackf77_dlarnv( &ione, iseed, &size, hC );
lapackf77_dlacpy( "Full", &M, &N, hC, &ldc, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb, queues0 );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_dsetmatrix( M, N, hB, lda, dB[dev], ldda, opts.queue );
// since when offset != 0, the GPU that does beta*C may not be 0,
// send initial hC to all GPUs.
magma_dsetmatrix( M, N, hC, lda, dC[dev], ldda, opts.queue );
}
trace_init( 1, opts.ngpu, nqueue, (magma_queue_t*) queues );
gpu_time = magma_sync_wtime(0);
magmablas_dsymm_mgpu(
MagmaLeft, MagmaLower, msize, N,
alpha, dA, ldda, offset,
dB, ldda,
beta, dC, ldda, dwork, dworksiz,
opts.ngpu, nb, queues, nqueue, events, nevents, gnode, ncmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "dsymm-m%d-n%d-nb%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magma_dsetmatrix( M, M, hA, lda, dA2, ldda, opts.queue );
magma_dsetmatrix( M, N, hB, lda, dB[0], ldda, opts.queue );
magma_dsetmatrix( M, N, hC, lda, dwork[0], ldda, opts.queue );
gpu_time2 = magma_sync_wtime(0);
magma_dsymm(
MagmaLeft, MagmaLower, msize, N,
alpha, dA2 + offset + offset*ldda, ldda,
dB[0], ldda,
beta, dwork[0], ldda, opts.queue );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||B||
Anorm = lapackf77_dlange("fro", &msize, &msize, hA + offset + offset*lda, &lda, work );
Anorm *= lapackf77_dlange("fro", &msize, &N, hB, &lda, work );
//printf( "A =" ); magma_dprint( M, M, hA, lda );
//printf( "B =" ); magma_dprint( M, N, hB, lda );
//printf( "C =" ); magma_dprint( M, N, hC, lda );
cpu_time = magma_wtime();
blasf77_dsymm( "Left", "Lower", &msize, &N,
&alpha, hA + offset + offset*lda, &lda,
hB, &lda,
&beta, hC, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
for (dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_dgetmatrix( M, N, dC[dev], ldda, hR, lda, opts.queue );
// compute relative error ||R||/||A||*||B||, where R := C_magma - C_lapack = R - C
size = ldc*N;
blasf77_daxpy( &size, &c_neg_one, hC, &ione, hR, &ione );
error = lapackf77_dlange("fro", &msize, &N, hR, &lda, work) / Anorm;
//printf( "R =" ); magma_dprint( M, N, hR, lda );
bool okay = (error < tol);
status += ! okay;
if (dev == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (okay ? "ok" : "failed") );
}
else {
printf( " dev %d %74s %8.2e %s\n", dev, "",
error, (okay ? "ok" : "failed") );
}
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hB );
TESTING_FREE_CPU( hC );
TESTING_FREE_PIN( hR );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
TESTING_FREE_DEV( dA[dev] );
TESTING_FREE_DEV( dB[dev] );
TESTING_FREE_DEV( dC[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( i = 0; i < nqueue; ++i ) {
magma_queue_destroy( queues[dev][i] );
}
for( i = 0; i < nevents; ++i ) {
magma_event_destroy( events[dev][i] );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
