int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
double *A, *B, *C, *C2, *LU;
double *dA, *dB, *dC1, *dC2;
double alpha = MAGMA_D_MAKE( 0.5, 0.1 );
double beta = MAGMA_D_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_err_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
printf("=========================================================================\n");
printf( "M %d, N %d, K %d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_dmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_dmalloc( &dA, size ); assert( err == 0 );
err = magma_dmalloc( &dB, size ); assert( err == 0 );
err = magma_dmalloc( &dC1, size ); assert( err == 0 );
err = magma_dmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_dlarnv( &ione, ISEED, &size, B );
lapackf77_dlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test DSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetmatrix( m, n, A, ld, dB, ld );
magma_dswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_dswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasDaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_dgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "dswap diff %.2g\n", error );
// ----- test IDAMAX
// get argmax of column of A
magma_dsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_idamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIdamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "idamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test DGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetvector( maxn, B, 1, dB, 1 );
magma_dsetvector( maxn, C, 1, dC1, 1 );
magma_dsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == 'N' ? m : n);
cublasDaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMV( m, n ) / 1e9;
printf( "dgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DSYMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetvector( m, B, 1, dB, 1 );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMV( m ) / 1e9;
printf( "dsymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_dlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_dgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_dsetmatrix( m, m, LU, ld, dA, ld );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "dtrsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test DGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == 'N');
bool ntb = (trans[ib] == 'N');
magma_dsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_dsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMM( m, n, k ) / 1e9;
printf( "dgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetmatrix( m, n, B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMM( side[is], m, n ) / 1e9;
printf( "dsymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYRK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_dsetmatrix( n, k, A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYRK( k, n ) / 1e9;
printf( "dsyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYR2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == 'N');
magma_dsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYR2K( k, n ) / 1e9;
printf( "dsyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRMM( side[is], m, n ) / 1e9;
printf( "dtrmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test DTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( side[is], m, n ) / 1e9;
printf( "dtrsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
return 0;
}
/***************************************************************************//**
Purpose
-------
DLAUUM computes the product U * U^H or L^H * L, where the triangular
factor U or L is stored in the upper or lower triangular part of
the array A.
If UPLO = MagmaUpper then the upper triangle of the result is stored,
overwriting the factor U in A.
If UPLO = MagmaLower then the lower triangle of the result is stored,
overwriting the factor L in A.
This is the blocked form of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the triangular factor stored in the array A
is upper or lower triangular:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the triangular factor U or L. N >= 0.
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = MagmaUpper, the upper triangle of A is
overwritten with the upper triangle of the product U * U^H;
if UPLO = MagmaLower, the lower triangle of A is overwritten with
the lower triangle of the product L^H * L.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -k, the k-th argument had an illegal value
@ingroup magma_lauum
*******************************************************************************/
extern "C" magma_int_t
magma_dlauum(
magma_uplo_t uplo, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i_, j_) ( A + (i_) + (j_)*lda )
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_) + (j_)*ldda)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#endif
/* Constants */
const double c_one = MAGMA_D_ONE;
const double d_one = MAGMA_D_ONE;
const char* uplo_ = lapack_uplo_const( uplo );
/* Local variables */
magma_int_t i, ib, ldda, nb;
magmaDouble_ptr dA;
bool upper = (uplo == MagmaUpper);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,n))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if (n == 0)
return *info;
nb = magma_get_dpotrf_nb( n );
ldda = magma_roundup( n, 32 );
if (MAGMA_SUCCESS != magma_dmalloc( &dA, n*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if (nb <= 1 || nb >= n) {
lapackf77_dlauum( uplo_, &n, A, &lda, info );
}
else if (upper) {
/* Compute the product U * U^H. */
// Computing 2nd block column (diagonal & above):
// [ u11 u12 u13 ] [ u11^H ] [ ... u12*u22^H + u13*u23^H ... ]
// [ u22 u23 ] * [ u12^H u22^H ] = [ ... u22*u22^H + u23*u23^H ... ]
// [ u33 ] [ u13^H u23^H u33^H ] [ ... ... ... ]
for (i=0; i < n; i += nb) {
ib = min( nb, n-i );
// Send diagonl block, u22
// This must finish before lauum below
magma_dsetmatrix( ib, ib,
A(i,i), lda,
dA(i,i), ldda, queues[0] );
// Send right of diagonl block, u23
magma_dsetmatrix_async( ib, n-i-ib,
A(i,i+ib), lda,
dA(i,i+ib), ldda, queues[1] );
// u12 = u12 * u22^H
magma_dtrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib, c_one,
dA(i,i), ldda,
dA(0,i), ldda, queues[0] );
// u22 = u22 * u22^H
lapackf77_dlauum( MagmaUpperStr, &ib, A(i,i), &lda, info );
magma_dsetmatrix_async( ib, ib,
A(i,i), lda,
dA(i,i), ldda, queues[0] );
if (i+ib < n) {
// wait for u23
magma_queue_sync( queues[1] );
// u12 += u13 * u23^H
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
i, ib, n-i-ib,
c_one, dA(0,i+ib), ldda,
dA(i,i+ib), ldda,
c_one, dA(0,i), ldda, queues[0] );
// u22 += u23 * u23^H
magma_dsyrk( MagmaUpper, MagmaNoTrans, ib, n-i-ib,
d_one, dA(i,i+ib), ldda,
d_one, dA(i,i), ldda, queues[0] );
}
// Get diagonal block & above of current column from device
// This could be on a different queue -- not needed until return
magma_dgetmatrix_async( i+ib, ib,
dA(0,i), ldda,
A(0,i), lda, queues[0] );
}
}
else {
/* Compute the product L^H * L. */
for (i=0; i < n; i += nb) {
ib = min( nb, n-i );
magma_dsetmatrix( ib, ib,
A(i,i), lda,
dA(i,i), ldda, queues[0] );
magma_dsetmatrix_async( n-i-ib, ib,
A(i+ib,i), lda,
dA(i+ib,i), ldda, queues[1] );
magma_dtrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib, i, c_one,
dA(i,i), ldda,
dA(i,0), ldda, queues[0] );
lapackf77_dlauum( MagmaLowerStr, &ib, A(i,i), &lda, info );
magma_dsetmatrix_async( ib, ib,
A(i,i), lda,
dA(i,i), ldda, queues[0] );
if (i+ib < n) {
magma_queue_sync( queues[1] );
magma_dgemm( MagmaConjTrans, MagmaNoTrans,
ib, i, n-i-ib,
c_one, dA(i+ib,i), ldda,
dA(i+ib,0), ldda,
c_one, dA(i,0), ldda, queues[0] );
magma_dsyrk( MagmaLower, MagmaConjTrans, ib, n-i-ib,
d_one, dA(i+ib,i), ldda,
d_one, dA(i,i), ldda, queues[0] );
}
magma_dgetmatrix_async( ib, i+ib,
dA(i,0), ldda,
A(i,0), lda, queues[0] );
}
}
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
}
/**
Purpose
-------
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_dposv_comp
********************************************************************/
extern "C" magma_int_t
magma_dpotrf_gpu(magma_uplo_t uplo, magma_int_t n,
double *dA, magma_int_t ldda, magma_int_t *info)
{
#define dA(i, j) (dA + (j)*ldda + (i))
magma_int_t j, jb, nb;
const char* uplo_ = lapack_uplo_const( uplo );
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *work;
double d_one = 1.0;
double d_neg_one = -1.0;
int upper = (uplo == MagmaUpper);
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_dpotrf_nb(n);
if (MAGMA_SUCCESS != magma_dmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
if ((nb <= 1) || (nb >= n)) {
/* Use unblocked code. */
magma_dgetmatrix_async( n, n, dA, ldda, work, n, stream[1] );
magma_queue_sync( stream[1] );
lapackf77_dpotrf(uplo_, &n, work, &n, info);
magma_dsetmatrix_async( n, n, work, n, dA, ldda, stream[1] );
}
else {
/* Use blocked code. */
if (upper) {
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j < n; j += nb) {
/* Update and factorize the current diagonal block and test
for non-positive-definiteness. Computing MIN */
jb = min(nb, (n-j));
magma_dsyrk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[0] );
if ( (j+jb) < n) {
/* Compute the current block row. */
magma_dgemm(MagmaConjTrans, MagmaNoTrans,
jb, (n-j-jb), j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_dpotrf(MagmaUpperStr, &jb, work, &jb, info);
magma_dsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[1] );
if (*info != 0) {
*info = *info + j;
break;
}
if ( (j+jb) < n) {
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, (n-j-jb),
c_one, dA(j, j ), ldda,
dA(j, j+jb), ldda);
}
}
}
else {
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j < n; j += nb) {
// Update and factorize the current diagonal block and test
// for non-positive-definiteness. Computing MIN
jb = min(nb, (n-j));
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[0] );
if ( (j+jb) < n) {
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
(n-j-jb), jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_dpotrf(MagmaLowerStr, &jb, work, &jb, info);
magma_dsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[1] );
if (*info != 0) {
*info = *info + j;
break;
}
if ( (j+jb) < n) {
magma_dtrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda);
}
}
}
}
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_dpotrf_gpu */
/**
Purpose
-------
DPOTRF_OOC computes the Cholesky factorization of a real symmetric
positive definite matrix A. This version does not require work
space on the GPU passed as input. GPU memory is allocated in the
routine. The matrix A may exceed the GPU memory.
The factorization has the form
A = U**H * U, if UPLO = MagmaUpper, or
A = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U**H * U or A = L * L**H.
\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,N).
@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.
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_dposv_comp
********************************************************************/
extern "C" magma_int_t
rr_dpotrf_m(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i, j) ( A + (j)*lda + (i))
#define dA(d, i, j) (dwork[(d)] + (j)*lddla + (i))
#define dT(d, i, j) ( dt[(d)] + (j)*ldda + (i))
#define dAup(d, i, j) (dwork[(d)] + (j)*NB + (i))
#define dTup(d, i, j) ( dt[(d)] + (j)*nb + (i))
/* Local variables */
double d_one = 1.0;
double d_neg_one = -1.0;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
const char* uplo_ = lapack_uplo_const( uplo );
int upper = (uplo == MagmaUpper);
double *dwork[MagmaMaxGPUs], *dt[MagmaMaxGPUs];
magma_int_t ldda, lddla, nb, iinfo, n_local[MagmaMaxGPUs], J2, d, ngpu0 = ngpu;
magma_int_t j, jj, jb, J, JB, NB, h;
magma_queue_t stream[MagmaMaxGPUs][3];
magma_event_t event[MagmaMaxGPUs][5];
magma_timer_t time_total=0, time_sum=0, time=0;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
nb = magma_get_dpotrf_nb(n);
if ( ngpu0 > n/nb ) {
ngpu = n/nb;
if ( n%nb != 0 ) ngpu ++;
} else {
ngpu = ngpu0;
}
//ldda = ((n+31)/32)*32;
ldda = ((n+nb-1)/nb)*nb;
lddla = ((nb*((n+nb*ngpu-1)/(nb*ngpu))+31)/32)*32;
/* figure out NB */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(double);
//MB = n; /* number of rows in the big panel */
NB = (magma_int_t)((0.8*freeMem - max(2,ngpu)*nb*ldda - (n+nb)*nb)/lddla); /* number of columns in the big panel */
//NB = min(5*nb,n);
if ( NB >= n ) {
#ifdef CHECK_DPOTRF_OOC
printf( " * still fits in GPU memory.\n" );
#endif
NB = n;
} else {
#ifdef CHECK_DPOTRF_OOC
printf( " * doesn't fit in GPU memory.\n" );
#endif
NB = (NB/nb) * nb; /* making sure it's devisable by nb */
}
#ifdef CHECK_DPOTRF_OOC
if ( NB != n ) printf( " * running in out-core mode (n=%d, NB=%d, nb=%d, lddla=%d, freeMem=%.2e).\n", n, NB, nb, lddla, (double)freeMem );
else printf( " * running in in-core mode (n=%d, NB=%d, nb=%d, lddla=%d, freeMem=%.2e).\n", n, NB, nb, lddla, (double)freeMem );
fflush(stdout);
#endif
for (d=0; d < ngpu; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_dmalloc( &dt[d], NB*lddla + max(2,ngpu)*nb*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork[d] = &dt[d][max(2,ngpu)*nb*ldda];
for( j=0; j < 3; j++ )
magma_queue_create( &stream[d][j] );
for( j=0; j < 5; j++ )
magma_event_create( &event[d][j] );
magma_device_sync(); // synch the device
}
magma_setdevice(0);
timer_start( time_total );
if (nb <= 1 || nb >= n) {
lapackf77_dpotrf(uplo_, &n, A, &lda, info);
} else {
/* Use hybrid blocked code. */
if (upper) {
/* =========================================================== *
* Compute the Cholesky factorization A = U'*U. *
* big panel is divided by block-row and distributed in block *
* column cyclic format */
/* for each big-panel */
for( J=0; J < n; J += NB ) {
JB = min(NB,n-J);
if ( ngpu0 > (n-J)/nb ) {
ngpu = (n-J)/nb;
if ( (n-J)%nb != 0 ) ngpu ++;
} else {
ngpu = ngpu0;
}
/* load the new big-panel by block-rows */
magma_dhtodpo( ngpu, uplo, JB, n, J, J, nb, A, lda, dwork, NB, stream, &iinfo);
/* update with the previous big-panels */
timer_start( time );
for( j=0; j < J; j += nb ) {
/* upload the diagonal of the block column (broadcast to all GPUs) */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_dsetmatrix_async( nb, JB,
A(j, J), lda,
dTup(d, 0, J), nb,
stream[d][0] );
n_local[d] = 0;
}
/* distribute off-diagonal blocks to GPUs */
for( jj=J+JB; jj < n; jj += nb ) {
d = ((jj-J)/nb)%ngpu;
magma_setdevice(d);
jb = min(nb, n-jj);
magma_dsetmatrix_async( nb, jb,
A(j, jj), lda,
dTup(d, 0, J+JB+n_local[d]), nb,
stream[d][0] );
n_local[d] += jb;
}
/* wait for the communication */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
}
/* update the current big-panel using the previous block-row */
/* -- process the big diagonal block of the big panel */
for( jj=0; jj < JB; jj += nb ) { // jj is 'local' column index within the big panel
d = (jj/nb)%ngpu;
J2 = jj/(nb*ngpu);
magma_setdevice(d);
magmablasSetKernelStream(stream[d][J2%2]); // the last stream (2) used to process off-diagonal
J2 = nb*J2;
jb = min(nb,JB-jj); // number of columns in this current block-row
magma_dgemm( MagmaConjTrans, MagmaNoTrans,
jj, jb, nb,
c_neg_one, dTup(d, 0, J ), nb,
dTup(d, 0, J+jj), nb,
c_one, dAup(d, 0, J2), NB);
magma_dsyrk(MagmaUpper, MagmaConjTrans, jb, nb,
d_neg_one, dTup(d, 0, J+jj), nb,
d_one, dAup(d, jj, J2), NB);
}
/* -- process the remaining big off-diagonal block of the big panel */
if ( n > J+JB ) {
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][2]);
/* local number of columns in the big panel */
n_local[d] = ((n-J)/(nb*ngpu))*nb;
if (d < ((n-J)/nb)%ngpu)
n_local[d] += nb;
else if (d == ((n-J)/nb)%ngpu)
n_local[d] += (n-J)%nb;
/* subtracting the local number of columns in the diagonal */
J2 = nb*(JB/(nb*ngpu));
if ( d < (JB/nb)%ngpu )
J2 += nb;
n_local[d] -= J2;
magma_dgemm( MagmaConjTrans, MagmaNoTrans,
JB, n_local[d], nb,
c_neg_one, dTup(d, 0, J ), nb,
dTup(d, 0, J+JB), nb,
c_one, dAup(d, 0, J2), NB);
}
}
/* wait for the previous updates */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
for( jj=0; jj < 3; jj++ )
magma_queue_sync( stream[d][jj] );
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
} /* end of updates with previous rows */
/* factor the big panel */
h = (JB+nb-1)/nb; // big diagonal of big panel will be on CPU
// using three streams
rr_dpotrf3_mgpu(ngpu, uplo, JB, n-J, J, J, nb,
dwork, NB, dt, ldda, A, lda, h, stream, event, &iinfo);
if ( iinfo != 0 ) {
*info = J+iinfo;
break;
}
time_sum += timer_stop( time );
/* upload the off-diagonal (and diagonal!!!) big panel */
magma_ddtohpo(ngpu, uplo, JB, n, J, J, nb, NB, A, lda, dwork, NB, stream, &iinfo);
//magma_ddtohpo(ngpu, uplo, JB, n, J, J, nb, 0, A, lda, dwork, NB, stream, &iinfo);
}
} else {
/* ========================================================= *
* Compute the Cholesky factorization A = L*L'. */
/* for each big-panel */
for( J=0; J < n; J += NB ) {
JB = min(NB,n-J);
if ( ngpu0 > (n-J)/nb ) {
ngpu = (n-J)/nb;
if ( (n-J)%nb != 0 ) ngpu ++;
} else {
ngpu = ngpu0;
}
/* load the new big-panel by block-columns */
magma_dhtodpo( ngpu, uplo, n, JB, J, J, nb, A, lda, dwork, lddla, stream, &iinfo);
/* update with the previous big-panels */
timer_start( time );
for( j=0; j < J; j += nb ) {
/* upload the diagonal of big panel */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_dsetmatrix_async( JB, nb,
A(J, j), lda,
dT(d, J, 0), ldda,
stream[d][0] );
n_local[d] = 0;
}
/* upload off-diagonals */
for( jj=J+JB; jj < n; jj += nb ) {
d = ((jj-J)/nb)%ngpu;
magma_setdevice(d);
jb = min(nb, n-jj);
magma_dsetmatrix_async( jb, nb,
A(jj, j), lda,
dT(d, J+JB+n_local[d], 0), ldda,
stream[d][0] );
n_local[d] += jb;
}
/* wait for the communication */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
}
/* update the current big-panel using the previous block-row */
for( jj=0; jj < JB; jj += nb ) { /* diagonal */
d = (jj/nb)%ngpu;
J2 = jj/(nb*ngpu);
magma_setdevice(d);
magmablasSetKernelStream(stream[d][J2%2]);
J2 = nb*J2;
jb = min(nb,JB-jj);
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
jb, jj, nb,
c_neg_one, dT(d, J+jj, 0), ldda,
dT(d, J, 0), ldda,
c_one, dA(d, J2, 0), lddla);
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, nb,
d_neg_one, dT(d, J+jj, 0), ldda,
d_one, dA(d, J2, jj), lddla);
}
if ( n > J+JB ) { /* off-diagonal */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][2]);
/* local number of columns in the big panel */
n_local[d] = (((n-J)/nb)/ngpu)*nb;
if (d < ((n-J)/nb)%ngpu)
n_local[d] += nb;
else if (d == ((n-J)/nb)%ngpu)
n_local[d] += (n-J)%nb;
/* subtracting local number of columns in diagonal */
J2 = nb*(JB/(nb*ngpu));
if ( d < (JB/nb)%ngpu )
J2 += nb;
n_local[d] -= J2;
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
n_local[d], JB, nb,
c_neg_one, dT(d, J+JB, 0), ldda,
dT(d, J, 0), ldda,
c_one, dA(d, J2, 0), lddla);
}
}
/* wait for the previous updates */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
for( jj=0; jj < 3; jj++ )
magma_queue_sync( stream[d][jj] );
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
}
/* factor the big panel */
h = (JB+nb-1)/nb; // big diagonal of big panel will be on CPU
// using three streams
magma_dpotrf3_mgpu(ngpu, uplo, n-J, JB, J, J, nb,
dwork, lddla, dt, ldda, A, lda, h, stream, event, &iinfo);
if ( iinfo != 0 ) {
*info = J+iinfo;
break;
}
time_sum += timer_stop( time );
/* upload the off-diagonal big panel */
magma_ddtohpo( ngpu, uplo, n, JB, J, J, nb, JB, A, lda, dwork, lddla, stream, &iinfo);
} /* end of for J */
} /* if upper */
} /* if nb */
timer_stop( time_total );
if ( ngpu0 > n/nb ) {
ngpu = n/nb;
if ( n%nb != 0 ) ngpu ++;
} else {
ngpu = ngpu0;
}
for (d=0; d < ngpu; d++ ) {
magma_setdevice(d);
for( j=0; j < 3; j++ ) {
magma_queue_destroy( stream[d][j] );
}
magma_free( dt[d] );
for( j=0; j < 5; j++ ) {
magma_event_destroy( event[d][j] );
}
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
timer_printf( "\n n=%d NB=%d nb=%d\n", (int) n, (int) NB, (int) nb );
timer_printf( " Without memory allocation: %f / %f = %f GFlop/s\n",
FLOPS_DPOTRF(n) / 1e9, time_total,
FLOPS_DPOTRF(n) / 1e9 / time_total );
timer_printf( " Performance %f / %f = %f GFlop/s\n",
FLOPS_DPOTRF(n) / 1e9, time_sum,
FLOPS_DPOTRF(n) / 1e9 / time_sum );
return *info;
} /* magma_dpotrf_ooc */
extern "C" magma_int_t
magma_dpotrf_gpu(char uplo, magma_int_t n,
double *dA, magma_int_t ldda, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**T * U, if UPLO = 'U', or
dA = L * L**T, if UPLO = 'L',
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of dA is stored;
= 'L': Lower triangle of dA is stored.
N (input) INTEGER
The order of the matrix dA. N >= 0.
dA (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix dA. If UPLO = 'U', the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**T * U or dA = L * L**T.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
magma_int_t j, jb, nb;
char uplo_[2] = {uplo, 0};
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *work;
double d_one = 1.0;
double d_neg_one = -1.0;
int upper = lapackf77_lsame(uplo_, "U");
*info = 0;
if ( (! upper) && (! lapackf77_lsame(uplo_, "L")) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_dpotrf_nb(n);
if (MAGMA_SUCCESS != magma_dmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
cudaStream_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else
stream[1] = current_stream;
if ((nb <= 1) || (nb >= n)) {
/* Use unblocked code. */
magma_dgetmatrix_async( n, n, dA, ldda, work, n, stream[1] );
magma_queue_sync( stream[1] );
lapackf77_dpotrf(uplo_, &n, work, &n, info);
magma_dsetmatrix_async( n, n, work, n, dA, ldda, stream[1] );
}
else {
/* Use blocked code. */
if (upper) {
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j<n; j+=nb) {
/* Update and factorize the current diagonal block and test
for non-positive-definiteness. Computing MIN */
jb = min(nb, (n-j));
magma_dsyrk(MagmaUpper, MagmaTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[0] );
if ( (j+jb) < n) {
/* Compute the current block row. */
magma_dgemm(MagmaTrans, MagmaNoTrans,
jb, (n-j-jb), j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_dpotrf(MagmaUpperStr, &jb, work, &jb, info);
magma_dsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[1] );
if (*info != 0) {
*info = *info + j;
break;
}
if ( (j+jb) < n) {
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, (n-j-jb),
c_one, dA(j, j ), ldda,
dA(j, j+jb), ldda);
}
}
}
else {
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j<n; j+=nb) {
// Update and factorize the current diagonal block and test
// for non-positive-definiteness. Computing MIN
jb = min(nb, (n-j));
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[0] );
if ( (j+jb) < n) {
magma_dgemm( MagmaNoTrans, MagmaTrans,
(n-j-jb), jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_dpotrf(MagmaLowerStr, &jb, work, &jb, info);
magma_dsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[1] );
if (*info != 0) {
*info = *info + j;
break;
}
if ( (j+jb) < n) {
magma_dtrsm(MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda);
}
}
}
}
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
return *info;
} /* magma_dpotrf_gpu */
/**
Purpose
-------
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix A. This version does not require work
space on the GPU passed as input. GPU memory is allocated in the
routine.
The factorization has the form
A = U**H * U, if uplo = MagmaUpper, or
A = L * L**H, if uplo = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
This uses multiple queues to overlap communication and computation.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If uplo = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If uplo = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U**H * U or A = L * L**H.
\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,N).
@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.
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_dposv_comp
********************************************************************/
extern "C" magma_int_t
magma_dpotrf(
magma_uplo_t uplo, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *info )
{
#define A(i_, j_) (A + (i_) + (j_)*lda)
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_) + (j_)*ldda)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#endif
/* Constants */
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const double d_one = 1.0;
const double d_neg_one = -1.0;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
bool upper = (uplo == MagmaUpper);
magma_int_t j, jb, ldda, nb;
magmaDouble_ptr dA = NULL;
/* Check arguments */
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
nb = magma_get_dpotrf_nb( n );
if (nb <= 1 || nb >= n) {
lapackf77_dpotrf( uplo_, &n, A, &lda, info );
}
else {
/* Use hybrid blocked code. */
ldda = magma_roundup( n, 32 );
magma_int_t ngpu = magma_num_gpus();
if ( ngpu > 1 ) {
/* call multi-GPU non-GPU-resident interface */
return magma_dpotrf_m( ngpu, uplo, n, A, lda, info );
}
if (MAGMA_SUCCESS != magma_dmalloc( &dA, n*ldda )) {
/* alloc failed so call the non-GPU-resident version */
return magma_dpotrf_m( ngpu, uplo, n, A, lda, info );
}
magma_queue_t queues[2] = { NULL, NULL };
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if (upper) {
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j < n; j += nb) {
/* Update and factorize the current diagonal block and test
for non-positive-definiteness. */
jb = min( nb, n-j );
magma_dsetmatrix_async( jb, n-j,
A(j, j), lda,
dA(j, j), ldda, queues[1] );
magma_dsyrk( MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, queues[0] );
if (j+jb < n) {
magma_dgemm( MagmaConjTrans, MagmaNoTrans,
jb, n-j-jb, j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda, queues[1] );
}
magma_queue_sync( queues[0] );
// this could be on any queue; it isn't needed until exit.
magma_dgetmatrix_async( j, jb,
dA(0, j), ldda,
A(0, j), lda, queues[0] );
lapackf77_dpotrf( MagmaUpperStr, &jb, A(j, j), &lda, info );
if (*info != 0) {
*info = *info + j;
break;
}
magma_dsetmatrix_async( jb, jb,
A(j, j), lda,
dA(j, j), ldda, queues[0] );
magma_queue_sync( queues[0] );
if (j+jb < n) {
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, n-j-jb,
c_one, dA(j, j ), ldda,
dA(j, j+jb), ldda, queues[1] );
}
}
}
else {
//used for timing CPU and GPU
int iter = 0;
float cpu_time = 0.0;
float gpu_time = 0.0;
double gpu_iter1_low = 2103.143311;
double gpu_iter1_high = 754.506104;
double cpu_iter1_low = 794.636108;
double cpu_iter1_high = 600.295227;
double gpu_pred_high = gpu_iter1_high;
double gpu_pred_low = gpu_iter1_low;
double cpu_pred_high = cpu_iter1_high;
double cpu_pred_low = cpu_iter1_low;
double ratio_split_freq = 0;
double time_until_interrupt = 0;
cudaEvent_t start_cpu, stop_cpu;
cudaEvent_t start_gpu, stop_gpu;
// switches for different modes
bool timing = false; //for initial setting only, greatly impact performance
bool dvfs = false; //turn on dvfs energy saving
bool relax = false; //turn on relax scheme
bool r2h = false; // turn on race to halt
//these parameters need to be tuned in future works.
double dvfs_converage = 0.5;
double prediction_offset_gpu = 0.65;
double prediction_offset_cpu = 0.65;
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j < n; j += nb) {
// Update and factorize the current diagonal block and test
// for non-positive-definiteness.
jb = min( nb, n-j );
magma_dsetmatrix_async( n-j, jb,
A(j, j), lda,
dA(j, j), ldda, queues[1] );
magma_dsyrk( MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_dgetmatrix_async( jb, jb,
dA(j,j), ldda,
A(j,j), lda, queues[0] );
if (timing) {
//start gpu timing
cudaEventCreate(&start_gpu);
cudaEventCreate(&stop_gpu);
cudaEventRecord(start_gpu, 0);
}
if (j+jb < n) {
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
n-j-jb, jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda, queues[1] );
}
double ratio_slack_pred = 1.0 - (double)nb/(n-iter*nb);
cpu_pred_high = cpu_pred_high * ratio_slack_pred;
cpu_pred_low = cpu_pred_low * ratio_slack_pred;
gpu_pred_high = gpu_pred_high * ratio_slack_pred * ratio_slack_pred;
gpu_pred_low = gpu_pred_low * ratio_slack_pred * ratio_slack_pred;
if (timing) {
printf("iter:%d GPU time pred:%f\n", iter, gpu_pred_high);
printf("iter:%d CPU time pred:%f\n", iter, cpu_pred_high);
}
if (iter < dvfs_converage*(n/nb)) {
if (cpu_pred_high > gpu_pred_high) { //slack on GPU
ratio_split_freq = (cpu_pred_high - gpu_pred_high) / (gpu_pred_high * ((gpu_iter1_low / gpu_iter1_high) - 1));
time_until_interrupt = gpu_pred_low * ratio_split_freq;
//printf("iter:%d time_until_interrupt:%f\n", iter, time_until_interrupt);
// printf("iter:%d ratio_split_freq:%f\n", iter, ratio_split_freq);
if (dvfs) {
if ((!relax) || (relax && ratio_split_freq > 0.05)) {
if (ratio_split_freq < 1)
dvfs_adjust(time_until_interrupt*prediction_offset_gpu, 'g');
else
dvfs_adjust(cpu_pred_high, 'g');
}
} else if (r2h) {
r2h_adjust(gpu_pred_high, cpu_pred_high - gpu_pred_high, 'g');
}
} else { //slack on CPU
ratio_split_freq = (gpu_pred_high - cpu_pred_high) / (cpu_pred_high * ((cpu_iter1_low / cpu_iter1_high) - 1));
time_until_interrupt = cpu_pred_low * ratio_split_freq;
if (dvfs) {
if ((!relax) || (relax && ratio_split_freq > 0.05)) {
if (ratio_split_freq < 1)
dvfs_adjust(time_until_interrupt*prediction_offset_cpu, 'c');
else
dvfs_adjust(gpu_pred_high, 'c');
}
} else if (r2h) {
r2h_adjust(cpu_pred_high, gpu_pred_high - cpu_pred_high, 'c');
}
}
}
if (timing) {
//end gpu timing
cudaEventRecord(stop_gpu, 0);
cudaEventSynchronize(stop_gpu);
cudaEventElapsedTime(&gpu_time, start_gpu, stop_gpu);
cudaEventDestroy(start_gpu);
cudaEventDestroy(stop_gpu);
//printf("iter:%d GPU time:%f\n", iter, gpu_time);
}
magma_queue_sync( queues[0] );
// this could be on any queue; it isn't needed until exit.
magma_dgetmatrix_async( jb, j,
dA(j, 0), ldda,
A(j, 0), lda, queues[0] );
if (timing) {
//start cpu timing
cudaEventCreate(&start_cpu);
cudaEventCreate(&stop_cpu);
cudaEventRecord(start_cpu, 0);
}
lapackf77_dpotrf( MagmaLowerStr, &jb, A(j, j), &lda, info );
if (timing) {
//end cpu timing
cudaEventRecord(stop_cpu, 0);
cudaEventSynchronize(stop_cpu);
cudaEventElapsedTime(&cpu_time, start_cpu, stop_cpu);
cudaEventDestroy(start_cpu);
cudaEventDestroy(stop_cpu);
// printf("iter:%d CPU time:%f\n", iter, cpu_time);
// if (gpu_time < cpu_time) {
// printf("slack: +\n");
// } else {
// printf("slack: -\n");
// }
}
if (*info != 0) {
*info = *info + j;
break;
}
magma_dsetmatrix_async( jb, jb,
A(j, j), lda,
dA(j, j), ldda, queues[0] );
magma_queue_sync( queues[0] );
if (j+jb < n) {
magma_dtrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-j-jb, jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda, queues[1] );
}
}
}
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
}
return *info;
} /* magma_dpotrf */
extern "C" magma_int_t
magma_dpotrf2_mgpu(
magma_int_t num_gpus, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb,
magmaDouble_ptr *d_lA, size_t d_lA_offset, magma_int_t ldda,
magmaDouble_ptr *d_lP, magma_int_t lddp,
double *a, magma_int_t lda, magma_int_t h,
magma_queue_t *queues,
magma_int_t *info )
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = 'U', or
dA = L * L**H, if UPLO = 'L',
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of dA is stored;
= 'L': Lower triangle of dA is stored.
N (input) INTEGER
The order of the matrix dA. N >= 0.
dA (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix dA. If UPLO = 'U', the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
magma_int_t j, jb, nb0, nb2, dd, d, id, j_local, j_local2, buf;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double d_one = 1.0;
double d_neg_one = -1.0;
magmaDouble_ptr dlpanel;
size_t dlpanel_offset;
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
magma_event_t events[MagmaMaxGPUs];
*info = 0;
if ( (uplo != MagmaUpper) && (uplo != MagmaLower) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if ((uplo != MagmaUpper) && num_gpus*ldda < max(1,n)) {
*info = -4;
} else if ((uplo == MagmaUpper) && ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
{
for( d=0; d<num_gpus; d++ ) {
/* local-n and local-ld */
if (uplo == MagmaUpper) {
n_local[d] = ((n/nb)/num_gpus)*nb;
if (d < (n/nb)%num_gpus)
n_local[d] += nb;
else if (d == (n/nb)%num_gpus)
n_local[d] += n%nb;
} else {
n_local[d] = ((m/nb)/num_gpus)*nb;
if (d < (m/nb)%num_gpus)
n_local[d] += nb;
else if (d == (m/nb)%num_gpus)
n_local[d] += m%nb;
}
}
/* Use blocked code. */
if (uplo == MagmaUpper)
{
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
for(j=0;j<m;j+=nb){
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (m-j));
if(j>0){
magma_queue_sync(queues[id*2]); // wait for the column on CPU
/* broadcast off-diagonal column to all gpus */
d = (j/nb+1)%num_gpus;
for(dd=0;dd<num_gpus;dd++){
if(d != id){
//magma_queue_sync(queues[2*d]);
magma_dsetmatrix_async( j, jb,
Aup(0,j), lda,
dlP(d,jb,0,buf), lddp,
queues[d*2], NULL );
}
d = (d+1)%num_gpus;
}
}
/* Update the current diagonal block */
if( j > 0 ) {
magma_dsyrk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dlA(id, 0, nb*j_local), ldda,
d_one, dlA(id, j, nb*j_local), ldda,
queues[2*id+1]);
}
/* send the diagonal to cpu */
magma_queue_sync(queues[2*id+1]);// wait for syrk
magma_dgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
queues[2*id], NULL);
if(j>0){
/* Compute the local block column of the panel. */
d = (j/nb+1)%num_gpus;
for(dd=0;dd<num_gpus;dd++){
j_local2 = j_local+1;
if(d>id) j_local2 --;
nb0 = nb*j_local2;
if(n_local[d]>nb0){
/* wait for the off-diagonal */
if(d!=id){
dlpanel = d_lP[d];
dlpanel_offset = dlP_offset(jb, 0, buf);
ldpanel = lddp;
/* wait for the offdiagonal column */
magma_queue_sync(queues[d*2]);
}else{
dlpanel = d_lA[d];
dlpanel_offset = dlA_offset(0, nb*j_local);
ldpanel = ldda;
}
/* update the panel */
magma_dgemm(MagmaConjTrans, MagmaNoTrans,
jb, n_local[d]-nb0, j,
c_neg_one, dlpanel, dlpanel_offset, ldpanel,
dlA(d, 0, nb0), ldda,
c_one, dlA(d, j, nb0), ldda,
queues[2*d+1]);
}
d = (d+1)%num_gpus;
}
}
/* factor the diagonal */
magma_queue_sync( queues[id*2] ); // wait for the diagonal
lapackf77_dpotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d == id ) {
dlpanel = d_lA[d];
dlpanel_offset = dlA_offset(j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = d_lP[d];
dlpanel_offset = dlP_offset(0, 0, buf);
ldpanel = lddp;
}
magma_dsetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, dlpanel_offset, ldpanel,
queues[d*2], NULL);
d = (d+1)%num_gpus;
}
} else {
magma_dsetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
queues[id*2], NULL );
}
/* panel-factorize the off-diagonal */
if((j+jb)<n){
d = (j/nb+1)%num_gpus;
for(dd=0;dd<num_gpus;dd++){
/* next column */
j_local2 = j_local+1;
if(d>id) j_local2--;
if( d == id ) {
dlpanel = d_lA[d];
dlpanel_offset = dlA_offset(j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = d_lP[d];
dlpanel_offset = dlP_offset(0, 0, buf);
ldpanel = lddp;
}
nb2 = n_local[d]-nb*j_local2;
nb0 = min(nb, nb2 );
magma_queue_sync( queues[2*d]); // wait for the diagonal
if(j+jb < m && d == (j/nb+1)%num_gpus){
/* owns the next column, look-ahead the column */
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2), ldda,
queues[2*d+1]);
/* send the column to cpu */
if(j+jb < m){
magma_queue_sync(queues[2*d+1]); // wait for lookahead
magma_dgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
queues[2*d], NULL);
}
/* update the remaining blocks */
nb2 = nb2 - nb0;
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
queues[2*d+1]);
}else if(nb2 > 0){
/* update the entire trailing matrix */
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2), ldda,
queues[d*2+1]);
}
d = (d+1)%num_gpus;
}
}
}
} else {
/* -------------------------------------------- */
/* Lower-triangular case */
/* Compute the Cholesky factorization A = L*L'. */
/* -------------------------------------------- */
for (j=0; j<n; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (n-j));
if( j > 0 ) {
/* needed on pluto... */
//magma_setdevice(id);
magma_queue_sync( queues[id*2] ); // wait for the column on CPU
/* broadcast offdiagonal row to all gpus */
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d != id ) {
/* wait for it on CPU */
//magma_queue_sync( queues[d*2] );
/* send it to GPU */
magma_dsetmatrix_async( jb, j,
Alo(j,0), lda,
dlPT(d,0,jb,buf), nb,
queues[d*2], NULL );
clFlush(queues[d*2]);
}
d = (d+1)%num_gpus;
}
}
/* Update the current diagonal block */
if( j > 0 ) {
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda,
queues[id*2+1]);
magma_queue_sync( queues[id*2+1] ); // wait for syrk
}
/* update the offdiagonal blocks */
if ( j > 0 ) {
/* compute the block-rows of the panel */
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
j_local2 = j_local+1;
if( d > id ) j_local2 --;
nb0 = nb*j_local2;
if( nb0 < n_local[d] ) {
if( d != id ) {
//dlpanel = dlPT(d);
dlpanel = d_lP[d];
dlpanel_offset = dlPT_offset(0, jb, buf);
ldpanel = nb;
/* wait for offdiagonal row */
magma_queue_sync(queues[d*2]);
} else {
dlpanel = d_lA[d];
dlpanel_offset = dlA_offset(nb*j_local, 0);
ldpanel = ldda;
}
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, dlpanel_offset, ldpanel,
c_one, dlA(d, nb0, j), ldda,
queues[d*2+1]);
}
d = (d+1)%num_gpus;
}
}
/* send the diagonal to cpu */
magma_dgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
queues[id*2], &events[id] );
clFlush(queues[id*2]);
/* factor the diagonal */
magma_queue_sync( queues[id*2] );
lapackf77_dpotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
if (*info != 0) {
printf("row number: %d\n", (int) j);
*info = *info + j;
break;
}
/* send the diagonal to gpus */
if ( (j+jb) < m ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d == id ) {
dlpanel = d_lA[d];
dlpanel_offset = dlA_offset(nb*j_local, j);
ldpanel = ldda;
} else {
//dlpanel = dlPT(d);
dlpanel = d_lP[d];
dlpanel_offset = dlPT_offset(0, 0, buf);
ldpanel = nb;
}
magma_dsetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, dlpanel_offset, ldpanel,
queues[d*2], NULL );
clFlush(queues[d*2]);
d = (d+1)%num_gpus;
}
} else {
magma_dsetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
queues[id*2], NULL );
clFlush(queues[id*2]);
}
/* factorize off-diagonal blocks */
if ( (j+jb) < m ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = d_lA[d];
dlpanel_offset = dlA_offset(nb*j_local, j);
ldpanel = ldda;
} else {
//dlpanel = dlPT(d);
dlpanel = d_lP[d];
dlpanel_offset = dlPT_offset(0, 0, buf);
ldpanel = nb;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2 );
magma_queue_sync(queues[d*2]);
// wait for the diagonal
if( j+jb < n && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead the column */
magma_dtrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2, j), ldda,
queues[d*2+1]);
/* send the column to cpu */
if( j+jb < n ) {
magma_queue_sync( queues[d*2+1] ); // wait for lookahead
magma_dgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
queues[d*2], NULL);
clFlush(queues[d*2]);
}
/* update the remaining blocks */
nb2 = nb2 - nb0;
magma_dtrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
queues[d*2+1]);
} else if( nb2 > 0 ) {
/* update the entire trailing matrix */
magma_dtrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2, j), ldda,
queues[d*2+1]);
}
d = (d+1)%num_gpus;
}
}
}
} /* end of else not upper */
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_queue_sync( queues[d*2] );
magma_queue_sync( queues[d*2+1] );
}
} /* end of not lapack */
return *info;
} /* magma_dpotrf_mgpu */
/**
Purpose
-------
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix A. This version does not require work
space on the GPU passed as input. GPU memory is allocated in the
routine.
The factorization has the form
A = U**H * U, if uplo = MagmaUpper, or
A = L * L**H, if uplo = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If uplo = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If uplo = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U**H * U or A = L * L**H.
\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,N).
@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.
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_dposv_comp
********************************************************************/
extern "C" magma_int_t
magma_dpotrf(
magma_uplo_t uplo, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i_, j_) (A + (j_)*lda + (i_))
#define dA(i_, j_) (dA + (j_)*ldda + (i_))
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ldda, nb;
magma_int_t j, jb;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magmaDouble_ptr dA;
double d_one = 1.0;
double d_neg_one = -1.0;
int upper = (uplo == MagmaUpper);
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
magma_int_t ngpu = magma_num_gpus();
if ( ngpu > 1 ) {
/* call multiple-GPU interface */
return magma_dpotrf_m(ngpu, uplo, n, A, lda, info);
}
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_dmalloc( &dA, (n)*ldda )) {
/* alloc failed so call the non-GPU-resident version */
return magma_dpotrf_m(ngpu, uplo, n, A, lda, info);
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[3];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
magma_queue_create( &stream[2] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
nb = magma_get_dpotrf_nb(n);
if (nb <= 1 || nb >= n) {
lapackf77_dpotrf(uplo_, &n, A, &lda, info);
} else {
/* Use hybrid blocked code. */
if (upper) {
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j < n; j += nb) {
/* Update and factorize the current diagonal block and test
for non-positive-definiteness. Computing MIN */
jb = min(nb, (n-j));
magma_dsetmatrix_async( jb, (n-j), A(j, j), lda, dA(j, j), ldda, stream[1]);
magma_dsyrk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, stream[0] );
if ( (j+jb) < n) {
magma_dgemm(MagmaConjTrans, MagmaNoTrans,
jb, (n-j-jb), j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda);
}
magma_dgetmatrix_async( j, jb,
dA(0, j), ldda,
A (0, j), lda, stream[2] );
magma_queue_sync( stream[0] );
lapackf77_dpotrf(MagmaUpperStr, &jb, A(j, j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
magma_dsetmatrix_async( jb, jb,
A(j, j), lda,
dA(j, j), ldda, stream[0] );
magma_queue_sync( stream[0] );
if ( (j+jb) < n ) {
magma_dtrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, (n-j-jb),
c_one, dA(j, j ), ldda,
dA(j, j+jb), ldda);
}
}
}
else {
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j < n; j += nb) {
// Update and factorize the current diagonal block and test
// for non-positive-definiteness. Computing MIN
jb = min(nb, (n-j));
magma_dsetmatrix_async( (n-j), jb, A(j, j), lda, dA(j, j), ldda, stream[1]);
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_dgetmatrix_async( jb, jb,
dA(j,j), ldda,
A(j,j), lda, stream[0] );
if ( (j+jb) < n) {
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
(n-j-jb), jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda);
}
magma_dgetmatrix_async( jb, j,
dA(j, 0), ldda,
A(j, 0), lda, stream[2] );
magma_queue_sync( stream[0] );
lapackf77_dpotrf(MagmaLowerStr, &jb, A(j, j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
magma_dsetmatrix_async( jb, jb,
A(j, j), lda,
dA(j, j), ldda, stream[0] );
magma_queue_sync( stream[0] );
if ( (j+jb) < n) {
magma_dtrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda);
}
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[2] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
magma_free( dA );
return *info;
} /* magma_dpotrf */
extern "C" magma_int_t
magma_dpotrf3_mgpu(int num_gpus, char uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb,
double **d_lA, magma_int_t ldda,
double **d_lP, magma_int_t lddp,
double *a, magma_int_t lda, magma_int_t h,
cudaStream_t stream[][3], cudaEvent_t event[][5],
magma_int_t *info )
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
Auxiliary subroutine for dpotrf2_ooc. It is multiple gpu interface to compute
Cholesky of a "rectangular" matrix.
The factorization has the form
dA = U**T * U, if UPLO = 'U', or
dA = L * L**T, if UPLO = 'L',
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of dA is stored;
= 'L': Lower triangle of dA is stored.
N (input) INTEGER
The order of the matrix dA. N >= 0.
dA (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix dA. If UPLO = 'U', the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**T * U or dA = L * L**T.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
magma_int_t j, jb, nb0, nb2, d, dd, id, j_local, j_local2, buf;
char uplo_[2] = {uplo, 0};
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double d_one = 1.0;
double d_neg_one = -1.0;
int upper = lapackf77_lsame(uplo_, "U");
double *dlpanel;
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
//cudaEvent_t event0[MagmaMaxGPUs], /* send row to CPU */
// event1[MagmaMaxGPUs], /* send diag to GPU */
// event2[MagmaMaxGPUs], /* offdiagonal update */
// event3[MagmaMaxGPUs], /* send row to GPU */
// event4[MagmaMaxGPUs]; /* lookahead */
const magma_int_t stream1 = 0, stream2 = 1, stream3 = 2;
double *d_dinvA[MagmaMaxGPUs][2], *d_x[MagmaMaxGPUs][2]; /* used by dtrsm_work */
*info = 0;
if ( (! upper) && (! lapackf77_lsame(uplo_, "L")) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (!upper && num_gpus*ldda < max(1,n)) {
*info = -4;
} else if (upper && ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* initialization */
for( d=0; d<num_gpus; d++ ) {
/* local-n and local-ld */
if (upper) {
n_local[d] = ((n/nb)/num_gpus)*nb;
if (d < (n/nb)%num_gpus)
n_local[d] += nb;
else if (d == (n/nb)%num_gpus)
n_local[d] += n%nb;
} else {
n_local[d] = ((m/nb)/num_gpus)*nb;
if (d < (m/nb)%num_gpus)
n_local[d] += nb;
else if (d == (m/nb)%num_gpus)
n_local[d] += m%nb;
}
//magma_setdevice(d);
//magma_event_create( &event0[d] );
//magma_event_create( &event1[d] );
//magma_event_create( &event2[d] );
//magma_event_create( &event3[d] );
//magma_event_create( &event4[d] );
}
/* == initialize the trace */
trace_init( 1, num_gpus, 3, (CUstream_st**)stream );
if (upper)
{
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
/* invert the diagonals
* Allocate device memory for the inversed diagonal blocks, size=m*NB
*/
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j<2; j++ ) {
cudaMalloc((void**)&d_dinvA[d][j], nb*nb*sizeof(double));
cudaMalloc((void**)&d_x[d][j], n*nb*sizeof(double));
cudaMemset(d_dinvA[d][j], 0, nb*nb*sizeof(double));
cudaMemset(d_x[d][j], 0, n*nb*sizeof(double));
}
}
magma_setdevice(0);
#endif
for (j=0; j<m; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (m-j));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
trace_gpu_start( id, stream1, "syrk", "syrk" );
magma_dsyrk(MagmaUpper, MagmaTrans, jb, j,
d_neg_one, dlA(id, 0, nb*j_local), ldda,
d_one, dlA(id, j, nb*j_local), ldda);
trace_gpu_end( id, stream1 );
}
/* send the diagonal to cpu on stream1 */
trace_gpu_start( id, stream1, "comm", "D to CPU" );
magma_dgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
stream[id][stream1] );
trace_gpu_end( id, stream1 );
/* update off-diagonal blocks in the panel */
if( j > 0 ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
j_local2 = j_local+1;
if( d > id ) j_local2 --;
nb0 = nb*j_local2;
if( n_local[d] > nb0 ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
if( d == id ) {
dlpanel = dlA(d, 0, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d, jb, 0, buf);
ldpanel = lddp;
magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
}
trace_gpu_start( d, stream2, "gemm", "gemm" );
magma_dgemm(MagmaTrans, MagmaNoTrans,
jb, n_local[d]-nb0, j,
c_neg_one, dlpanel, ldpanel,
dlA(d, 0, nb0), ldda,
c_one, dlA(d, j, nb0), ldda);
trace_gpu_end( d, stream2 );
magma_event_record( event[d][2], stream[d][stream2] );
}
d = (d+1)%num_gpus;
}
}
/* wait for panel and factorize it on cpu */
magma_setdevice(id);
magma_queue_sync( stream[id][stream1] );
trace_cpu_start( 0, "getrf", "getrf" );
lapackf77_dpotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus on stream1 */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d, 0, 0, buf);
ldpanel = lddp;
}
magma_setdevice(d);
trace_gpu_start( d, stream1, "comm", "comm" );
magma_dsetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, ldpanel,
stream[d][stream1] );
trace_gpu_end( d, stream1 );
magma_event_record( event[d][1], stream[d][stream1] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
trace_gpu_start( id, stream1, "comm", "comm" );
magma_dsetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
stream[id][stream1] );
trace_gpu_end( id, stream1 );
}
/* panel-factorize the off-diagonal */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = dlA(d,j,nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d, 0, 0, buf);
ldpanel = lddp;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2);
magma_setdevice(d);
//magma_queue_sync( stream[d][stream1] ); // synch on chol for remaining update
//magma_queue_sync( stream[d][stream2] );
if( j+jb < m && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead next block on stream1 */
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for gemm update
magmablasSetKernelStream(stream[d][stream1]);
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
magmablas_dtrsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d][0], d_x[d][0] );
#else
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
magma_event_record( event[d][4], stream[d][stream1] );
trace_gpu_end( d, stream1 );
} else if( nb2 > 0 ) {
/* update all the blocks on stream2 */
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for cholesky factor
trace_gpu_start( d, stream2, "trsm", "trsm" );
magmablasSetKernelStream(stream[d][stream2]);
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
magmablas_dtrsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream2 );
}
d = (d+1)%num_gpus;
} /* end of for */
/* ======================================================================================== */
d = (j/nb+1)%num_gpus;
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
nb0 = min(nb, n_local[d]-nb*j_local2 );
/* even on 1 gpu, off-diagonals are copied to cpu (synchronize at the end). *
* so we have the Cholesky factor, but only diagonal submatrix of the big panel, *
* on cpu at the end. */
if( j+jb < m ) {
int d2, id2, j2, buf2;
magma_setdevice(d);
/* make sure all the previous sets are done */
if( h < num_gpus ) {
/* > offdiagonal */
for( d2=0; d2<num_gpus; d2++ ) {
j2 = j - (1+d2)*nb;
if( j2 < 0 ) break;
id2 = (j2/nb)%num_gpus;
magma_queue_wait_event( stream[d][stream3], event[id2][0] );
}
/* > diagonal */
for( d2=0; d2<num_gpus; d2++ ) {
j2 = j - d2*nb;
if( j2 < 0 ) break;
id2 = (j2/nb)%num_gpus;
magma_queue_wait_event( stream[d][stream3], event[id2][1] );
}
}
/* lookahead */
magma_queue_wait_event( stream[d][stream3], event[d][4] );
trace_gpu_start( d, stream3, "comm", "row to CPU" );
magma_dgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
stream[d][stream3] );
trace_gpu_end( d, stream3 );
magma_event_record( event[d][3], stream[d][stream3] );
/* needed on pluto */
magma_queue_sync( stream[d][stream3] );
/* wait for the off-diagonal on cpu */
//magma_setdevice(id);
//magma_queue_sync( stream[id][stream3] );
/* broadcast rows to gpus on stream2 */
buf2 = ((j+jb)/nb)%num_gpus;
for( d2=0; d2<num_gpus; d2++ ) {
if( d2 != d )
{
magma_setdevice(d2);
trace_gpu_start( d2, stream3, "comm", "row to GPUs" );
magma_queue_wait_event( stream[d2][stream3], event[d][3] ); // rows arrived at cpu on stream3
magma_dsetmatrix_async( j+jb, nb0,
Aup(0,j+jb), lda,
dlP(d2,nb0,0,buf2), lddp,
stream[d2][stream3] );
trace_gpu_end( d2, stream3 );
magma_event_record( event[d2][0], stream[d2][stream3] );
}
}
}
/* ======================================================================================== */
/* gpu owning the next column */
/* after look ahead, update the remaining blocks */
if( j+jb < m ) /* no update on the last block column */
{
d = (j/nb+1)%num_gpus;
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d, 0, 0, buf);
ldpanel = lddp;
}
nb0 = min(nb, n_local[d]-nb*j_local2 );
nb2 = n_local[d]-nb*j_local2 - nb0;
/* update the remaining blocks */
if( nb2 > 0 ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for cholesky factor
trace_gpu_start( d, stream2, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
magmablas_dtrsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda);
#endif
trace_gpu_end( d, stream2 );
}
}
} /* end of dtrsm */
} /* end of for j=1, .., n */
} else {
/* ---------------------------------------------- */
/* Lower-triangular case */
/* > Compute the Cholesky factorization A = L*L'. */
/* ---------------------------------------------- */
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
/*
* Allocate device memory for the inversed diagonal blocks, size=N*BLOCK_SIZE
*/
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j<2; j++ ) {
cudaMalloc((void**)&d_dinvA[d][j], nb*nb*sizeof(double));
cudaMalloc((void**)&d_x[d][j], nb*m *sizeof(double));
cudaMemset(d_dinvA[d][j], 0, nb*nb*sizeof(double));
cudaMemset(d_x[d][j], 0, nb* m*sizeof(double));
}
}
magma_setdevice(0);
#endif
for (j=0; j<n; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (n-j));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda);
}
/* send the diagonal to cpu on stream1 */
magma_dgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
stream[id][stream1] );
/* update off-diagonal blocks of the panel */
if( j > 0 ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
j_local2 = j_local+1;
if( d > id ) j_local2 --;
nb0 = nb*j_local2;
if( nb0 < n_local[d] ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
if( d == id ) {
dlpanel = dlA(d, nb*j_local, 0);
ldpanel = ldda;
} else {
dlpanel = dlPT(d,0,jb,buf);
ldpanel = nb;
magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
}
magma_dgemm( MagmaNoTrans, MagmaTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, ldpanel,
c_one, dlA(d, nb0, j), ldda);
magma_event_record( event[d][2], stream[d][stream2] );
}
d = (d+1)%num_gpus;
}
}
/* wait for the panel and factorized it on cpu */
magma_setdevice(id);
magma_queue_sync( stream[id][stream1] );
lapackf77_dpotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus on stream1 */
if ( (j+jb) < m) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
magma_setdevice(d);
magma_dsetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, ldpanel,
stream[d][stream1] );
magma_event_record( event[d][1], stream[d][stream1] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
magma_dsetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
stream[id][stream1] );
}
/* panel factorize the off-diagonal */
if ( (j+jb) < m) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2 );
magma_setdevice(d);
if( j+nb < n && d == (j/nb+1)%num_gpus ) { /* owns next column, look-ahead next block on stream1 */
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for gemm update
magmablasSetKernelStream(stream[d][stream1]);
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
magmablas_dtrsm_work( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
d_dinvA[d][0], d_x[d][0] );
#else
magma_dtrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
magma_event_record( event[d][4], stream[d][stream1] );
} else if( nb2 > 0 ) { /* other gpus updating all the blocks on stream2 */
/* update the entire column */
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for the cholesky factor
magmablasSetKernelStream(stream[d][stream2]);
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
magmablas_dtrsm_work( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_dtrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
}
d = (d+1)%num_gpus;
} /* end for d */
/* ======================================================================================== */
d = (j/nb+1)%num_gpus;
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
nb0 = min(nb, n_local[d]-nb*j_local2 );
/* even on 1 gpu, we copy off-diagonal to cpu (but don't synchronize). */
/* so we have the Cholesky factor on cpu at the end. */
if( j+jb < n ) {
int d2, id2, j2, buf2;
magma_setdevice(d);
/* make sure all the previous sets are done */
if( h < num_gpus ) {
/* > offdiagonal */
for( d2=0; d2<num_gpus; d2++ ) {
j2 = j - (1+d2)*nb;
if( j2 < 0 ) break;
id2 = (j2/nb)%num_gpus;
magma_queue_wait_event( stream[d][stream3], event[id2][0] );
}
/* > diagonal */
for( d2=0; d2<num_gpus; d2++ ) {
j2 = j - d2*nb;
if( j2 < 0 ) break;
id2 = (j2/nb)%num_gpus;
magma_queue_wait_event( stream[d][stream3], event[id2][1] );
}
}
// lookahead done
magma_queue_wait_event( stream[d][stream3], event[d][4] );
magma_dgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
stream[d][stream3] );
magma_event_record( event[d][3], stream[d][stream3] );
/* syn on rows on CPU, seem to be needed on Pluto */
magma_queue_sync( stream[d][stream3] );
/* broadcast the rows to gpus */
buf2 = ((j+jb)/nb)%num_gpus;
for( d2=0; d2<num_gpus; d2++ ) {
if( d2 != d )
{
magma_setdevice(d2);
magma_queue_wait_event( stream[d2][stream3], event[d][3] ); // getmatrix done
magma_dsetmatrix_async( nb0, j+jb,
Alo(j+jb,0), lda,
dlPT(d2,0,nb0,buf2), nb,
stream[d2][stream3] );
magma_event_record( event[d2][0], stream[d2][stream3] );
}
}
}
/* ======================================================================================== */
/* gpu owing the next column updates remaining blocks on stream2 */
if( j+nb < n ) { // no lookahead on the last block column
d = (j/nb+1)%num_gpus;
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d,0,0,buf);
ldpanel = nb;
}
nb0 = min(nb, n_local[d]-nb*j_local2 );
nb2 = n_local[d] - j_local2*nb - nb0;
if( nb2 > 0 ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
/* update the remaining blocks in the column */
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
magmablas_dtrsm_work( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_dtrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda);
#endif
}
}
}
}
} /* end of else not upper */
/* == finalize the trace == */
trace_finalize( "dpotrf.svg","trace.css" );
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magma_queue_sync( stream[d][2] );
magmablasSetKernelStream(NULL);
//magma_event_destroy( event0[d] );
//magma_event_destroy( event1[d] );
//magma_event_destroy( event2[d] );
//magma_event_destroy( event3[d] );
//magma_event_destroy( event4[d] );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(DTRSM_WORK)
for( j=0; j<2; j++ ) {
magma_free( d_dinvA[d][j] );
magma_free( d_x[d][j] );
}
#endif
}
magma_setdevice(0);
return *info;
} /* magma_dpotrf_mgpu */
extern "C" magma_int_t
magma_dpotrf(magma_uplo_t uplo, magma_int_t n,
double *a, magma_int_t lda, magma_int_t *info,
magma_queue_t* queue )
{
/* -- clMAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix A. This version does not require work
space on the GPU passed as input. GPU memory is allocated in the
routine.
The factorization has the form
A = U**T * U, if UPLO = 'U', or
A = L * L**T, if UPLO = 'L',
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U**T * U or A = L * L**T.
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
INFO (output) 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.
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
magma_int_t ldda, nb, j, jb;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magmaDouble_ptr work;
double d_one = 1.0;
double d_neg_one = -1.0;
*info = 0;
if( (uplo != MagmaUpper) && (uplo != MagmaLower) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
magma_int_t num_gpus = magma_num_gpus();
if( num_gpus > 1 ) {
/* call multiple-GPU interface */
printf("multiple-GPU verison not implemented\n");
return MAGMA_ERR_NOT_IMPLEMENTED;
//return magma_dpotrf_m(num_gpus, uplo, n, a, lda, info);
}
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_dmalloc( &work, (n)*ldda )) {
/* alloc failed so call the non-GPU-resident version */
printf("non-GPU-resident version not implemented\n");
return MAGMA_ERR_NOT_IMPLEMENTED;
//return magma_dpotrf_m(num_gpus, uplo, n, a, lda, info);
}
nb = magma_get_dpotrf_nb(n);
if (nb <= 1 || nb >= n) {
lapackf77_dpotrf(lapack_uplo_const(uplo), &n, a, &lda, info);
} else {
/* Use hybrid blocked code. */
if (uplo == MagmaUpper) {
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j<n; j += nb) {
/* Update and factorize the current diagonal block and test
for non-positive-definiteness. Computing MIN */
jb = min(nb, (n-j));
magma_dsetmatrix_async( jb, (n-j), A(j, j), 0, lda, dA(j, j), ldda, queue[1], NULL);
magma_dsyrk(MagmaUpper, MagmaTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda, queue[1]);
magma_queue_sync( queue[1] );
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), 0, lda, queue[0], NULL );
if ( (j+jb) < n) {
magma_dgemm(MagmaTrans, MagmaNoTrans,
jb, (n-j-jb), j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda, queue[1]);
}
magma_queue_sync( queue[0] );
magma_dgetmatrix_async( j, jb,
dA(0, j), ldda,
A (0, j), 0, lda, queue[0], NULL );
lapackf77_dpotrf(MagmaUpperStr, &jb, A(j, j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
magma_dsetmatrix_async( jb, jb,
A(j, j), 0, lda,
dA(j, j), ldda, queue[0], NULL );
magma_queue_sync( queue[0] );
if ( (j+jb) < n ) {
magma_dtrsm(MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, (n-j-jb),
c_one, dA(j, j ), ldda,
dA(j, j+jb), ldda, queue[1] );
}
}
}
else {
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j<n; j+=nb) {
// Update and factorize the current diagonal block and test
// for non-positive-definiteness. Computing MIN
jb = min(nb, (n-j));
magma_dsetmatrix_async( (n-j), jb, A(j, j), 0, lda, dA(j, j), ldda, queue[1], NULL);
magma_dsyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda, queue[1]);
magma_queue_sync( queue[1] );
magma_dgetmatrix_async( jb, jb,
dA(j,j), ldda,
A(j,j), 0, lda, queue[0], NULL );
if ( (j+jb) < n) {
magma_dgemm( MagmaNoTrans, MagmaTrans,
(n-j-jb), jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda, queue[1]);
}
magma_queue_sync( queue[0] );
magma_dgetmatrix_async( jb, j,
dA(j, 0), ldda,
A(j, 0), 0, lda, queue[1], NULL );
lapackf77_dpotrf(MagmaLowerStr, &jb, A(j, j), &lda, info);
if (*info != 0){
*info = *info + j;
break;
}
magma_dsetmatrix_async( jb, jb,
A(j, j), 0, lda,
dA(j, j), ldda, queue[0], NULL );
magma_queue_sync( queue[0] );
if ( (j+jb) < n) {
magma_dtrsm(MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda, queue[1]);
}
}
}
}
magma_free( work );
return *info;
} /* magma_dpotrf */
extern "C" magma_int_t
magma_dlauum_gpu(char uplo, magma_int_t n,
double *dA, magma_int_t ldda, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DLAUUM computes the product U * U' or L' * L, where the triangular
factor U or L is stored in the upper or lower triangular part of
the array dA.
If UPLO = 'U' or 'u' then the upper triangle of the result is stored,
overwriting the factor U in dA.
If UPLO = 'L' or 'l' then the lower triangle of the result is stored,
overwriting the factor L in dA.
This is the blocked form of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the triangular factor stored in the array dA
is upper or lower triangular:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the triangular factor U or L. N >= 0.
dA (input/output) DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = 'U', the upper triangle of dA is
overwritten with the upper triangle of the product U * U';
if UPLO = 'L', the lower triangle of dA is overwritten with
the lower triangle of the product L' * L.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -k, the k-th argument had an illegal value
===================================================================== */
/* Local variables */
char uplo_[2] = {uplo, 0};
magma_int_t nb, i, ib;
double d_one = MAGMA_D_ONE;
double c_one = MAGMA_D_ONE;
double *work;
int upper = lapackf77_lsame(uplo_, "U");
*info = 0;
if ((! upper) && (! lapackf77_lsame(uplo_, "L")))
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,n))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_dpotrf_nb(n);
if (MAGMA_SUCCESS != magma_dmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n)
{
magma_dgetmatrix( n, n, dA, ldda, work, n );
lapackf77_dlauum(uplo_, &n, work, &n, info);
magma_dsetmatrix( n, n, work, n, dA, ldda );
}
else
{
if (upper)
{
/* Compute inverse of upper triangular matrix */
for (i=0; i < n; i += nb)
{
ib = min(nb, (n-i));
/* Compute the product U * U'. */
magma_dtrmm( MagmaRight, MagmaUpper,
MagmaTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
magma_dgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_dlauum(MagmaUpperStr, &ib, work, &ib, info);
magma_dsetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if(i+ib < n)
{
magma_dgemm( MagmaNoTrans, MagmaTrans,
i, ib, (n-i-ib), c_one, dA(0,i+ib),
ldda, dA(i, i+ib), ldda, c_one,
dA(0,i), ldda);
magma_dsyrk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
}
}
else
{
/* Compute the product L' * L. */
for(i=0; i<n; i=i+nb)
{
ib=min(nb,(n-i));
magma_dtrmm( MagmaLeft, MagmaLower,
MagmaTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
magma_dgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_dlauum(MagmaLowerStr, &ib, work, &ib, info);
magma_dsetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if((i+ib) < n)
{
magma_dgemm( MagmaTrans, MagmaNoTrans,
ib, i, (n-i-ib), c_one, dA( i+ib,i),
ldda, dA(i+ib, 0),ldda, c_one,
dA(i,0), ldda);
magma_dsyrk( MagmaLower, MagmaTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
}
/**
Purpose
-------
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
d_lA DOUBLE PRECISION array of pointers on the GPU, dimension (ngpu)
On entry, the symmetric matrix dA distributed over GPUs
(dl_A[d] points to the local matrix on the d-th GPU).
It is distributed in 1D block column or row cyclic (with the
block size of nb) if UPLO = MagmaUpper or MagmaLower, respectively.
If UPLO = MagmaUpper, the leading N-by-N upper triangular
part of dA contains the upper triangular part of the matrix dA,
and the strictly lower triangular part of dA is not referenced.
If UPLO = MagmaLower, the leading N-by-N lower triangular part
of dA contains the lower triangular part of the matrix dA, and
the strictly upper triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_dposv_comp
********************************************************************/
extern "C" magma_int_t
magma_dpotrf_mgpu_right(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n,
magmaDouble_ptr d_lA[], magma_int_t ldda,
magma_int_t *info )
{
#define dlA(id, i, j) (d_lA[(id)] + (j) * ldda + (i))
#define dlP(id, i, j) (d_lP[(id)] + (j) * ldda + (i))
#define panel(j) (panel + (j))
#define tmppanel(j) (tmppanel + (j))
#define tmpprevpanel(j) (tmpprevpanel + (j))
#define STREAM_ID(i) (nqueue > 1 ? 1+((i)/nb)%(nqueue-1) : 0)
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double d_one = 1.0;
double d_neg_one = -1.0;
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t j, nb, d, id, j_local, blkid, crosspoint, prevtrsmrows=0, nqueue = 5;
double *panel, *tmppanel0, *tmppanel1, *tmppanel, *tmpprevpanel;
double *d_lP[MagmaMaxGPUs], *dlpanel, *dlpanels[MagmaMaxGPUs];
magma_int_t rows, trsmrows, igpu, n_local[MagmaMaxGPUs], ldpanel;
magma_queue_t queues[MagmaMaxGPUs][10];
*info = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
nb = magma_get_dpotrf_nb(n);
ldpanel = ldda;
magma_setdevice(0);
if (MAGMA_SUCCESS != magma_dmalloc_pinned( &panel, 2 * nb * ldpanel )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
tmppanel0 = panel;
tmppanel1 = tmppanel0 + nb * ldpanel;
if ((nb <= 1) || (nb >= n)) {
// Use unblocked code.
magma_dgetmatrix( n, n, dlA(0, 0, 0), ldda, panel, ldpanel);
lapackf77_dpotrf( uplo_, &n, panel, &ldpanel, info);
magma_dsetmatrix( n, n, panel, ldpanel, dlA(0, 0, 0), ldda );
} else {
for( d = 0; d < ngpu; d++ ) {
// local-n and local-ld
n_local[d] = ((n / nb) / ngpu) * nb;
if (d < (n / nb) % ngpu)
n_local[d] += nb;
else if (d == (n / nb) % ngpu)
n_local[d] += n % nb;
magma_setdevice(d);
magma_device_sync();
if (MAGMA_SUCCESS != magma_dmalloc( &d_lP[d], nb * ldda )) {
for( j = 0; j < d; j++ ) {
magma_setdevice(j);
magma_free( d_lP[d] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
for( j=0; j < nqueue; j++ ) {
magma_queue_create( &queues[d][j] );
}
}
//#define ENABLE_TIMER
#if defined (ENABLE_TIMER)
real_Double_t therk[4], tmtc, tcchol, tctrsm, tctm, tmnp, tcnp;
real_Double_t ttot_herk[4] = {0,0,0,0}, ttot_mtc = 0, ttot_cchol = 0, ttot_ctrsm = 0, ttot_ctm = 0, ttot_mnp = 0, ttot_cnp = 0;
printf("\n\n %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s\n",
"j", "nb", "row", "mtc", "CPU_np", "panel", "ctrsm", "CH+TRSM", "CPU", "dsyrk[0]", "dsyrk[1]", "dsyrk[2]", "dsyrk[3]", "ctm P", "gpu_np");
printf(" ====================================================================================================\n");
#endif
// Use blocked code.
if (uplo == MagmaUpper) {
printf( " === not supported, yet ===\n" );
} else {
blkid = -1;
if (ngpu == 4)
crosspoint = n;
else if (ngpu == 3)
crosspoint = n;
else if (ngpu == 2)
crosspoint = 20160;
else
crosspoint = 0;
crosspoint = 0; //n; //n -- > gpu always does next panel, 0 --> cpu always does next panel
crosspoint = n;
#if defined (ENABLE_TIMER)
real_Double_t tget = magma_wtime(), tset = 0.0, ttot = 0.0;
#endif
if ( n > nb ) {
// send first panel to cpu
magma_setdevice(0);
tmppanel = tmppanel0;
magma_dgetmatrix_async(n, nb,
dlA(0, 0, 0), ldda,
tmppanel(0), ldpanel,
queues[0][0] );
}
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_device_sync();
}
tget = magma_wtime()-tget;
#endif
// Compute the Cholesky factorization A = L*L'
for (j = 0; (j + nb) < n; j += nb) {
#if defined (ENABLE_TIMER)
therk[0] = therk[1] = therk[2] = therk[3] = tmtc = tcchol = tctrsm = tctm = tmnp = tcnp = 0.0;
#endif
blkid += 1;
tmppanel = (blkid % 2 == 0) ? tmppanel0 : tmppanel1;
// Set the gpu number that holds the current panel
id = (j / nb) % ngpu;
magma_setdevice(id);
// Set the local index where the current panel is
j_local = j / (nb * ngpu) * nb;
rows = n - j;
// Wait for the panel on cpu
magma_queue_sync( queues[id][0] );
if (j > 0 && prevtrsmrows > crosspoint) {
#if defined (ENABLE_TIMER)
tcnp = magma_wtime();
#endif
tmpprevpanel = ((blkid - 1) % 2) == 0 ? tmppanel0 : tmppanel1;
blasf77_dgemm( MagmaNoTransStr, MagmaConjTransStr,
&rows, &nb, &nb,
&c_neg_one, tmpprevpanel(j), &ldpanel,
tmpprevpanel(j), &ldpanel,
&c_one, tmppanel(j), &ldpanel );
#if defined (ENABLE_TIMER)
tcnp = magma_wtime() - tcnp;
ttot_cnp += tcnp;
#endif
}
#if defined (ENABLE_TIMER)
tcchol = magma_wtime();
#endif
lapackf77_dpotrf(MagmaLowerStr, &nb, tmppanel(j), &ldpanel, info);
if (*info != 0) {
*info = *info + j;
break;
}
#if defined (ENABLE_TIMER)
tcchol = magma_wtime() - tcchol;
ttot_cchol += tcchol;
tctrsm = magma_wtime();
#endif
trsmrows = rows - nb;
if (trsmrows > 0) {
blasf77_dtrsm(MagmaRightStr, MagmaLowerStr, MagmaConjTransStr, MagmaNonUnitStr,
&trsmrows, &nb,
&c_one, tmppanel(j), &ldpanel,
tmppanel(j + nb), &ldpanel);
}
#if defined (ENABLE_TIMER)
tctrsm = magma_wtime() - tctrsm;
ttot_ctrsm += tctrsm;
tctm = magma_wtime();
#endif
d = (id + 1) % ngpu;
// send current panel to gpus
for (igpu = 0; igpu < ngpu; igpu++, d = (d + 1) % ngpu ) {
magma_int_t myrows = 0;
magma_int_t row_offset = 0;
if ( d == id ) {
dlpanel = dlA(d, j, j_local);
myrows = rows;
row_offset = 0;
} else {
dlpanel = dlP(d, 0, 0);
myrows = trsmrows;
row_offset = nb;
}
if (myrows > 0) {
magma_setdevice(d);
magma_dsetmatrix_async(myrows, nb,
tmppanel(j + row_offset), ldpanel,
dlpanel, ldda, queues[d][0] );
}
}
/* make sure panel is on GPUs */
d = (id + 1) % ngpu;
for (igpu = 0; igpu < ngpu; igpu++, d = (d + 1) % ngpu ) {
magma_setdevice(d);
magma_queue_sync( queues[d][0] );
}
#if defined (ENABLE_TIMER)
tctm = magma_wtime() - tctm;
ttot_ctm += tctm;
#endif
if ( (j + nb) < n) {
magma_int_t offset = 0;
magma_int_t row_offset = 0;
if (j + nb + nb < n) {
d = (id + 1) % ngpu;
magma_setdevice(d);
magma_int_t j_local2 = (j + nb) / (nb * ngpu) * nb;
if (trsmrows <= crosspoint) {
#if defined (ENABLE_TIMER)
tmnp = magma_wtime();
#endif
// do gemm on look ahead panel
if ( d == id ) {
dlpanel = dlA(d, j + nb, j_local);
} else {
dlpanel = dlP(d, 0, 0);
}
magmablasSetKernelStream( queues[d][STREAM_ID(j_local2)] );
#define DSYRK_ON_DIAG
#ifdef DSYRK_ON_DIAG
magma_dsyrk( MagmaLower, MagmaNoTrans,
nb, nb,
d_neg_one, dlpanel, ldda,
d_one, dlA(d, j + nb, j_local2), ldda);
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows-nb, nb, nb,
c_neg_one, dlpanel+nb, ldda,
dlpanel, ldda,
c_one, dlA(d, j + nb +nb, j_local2), ldda);
#else
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows, nb, nb,
c_neg_one, dlpanel, ldda,
dlpanel, ldda,
c_one, dlA(d, j + nb, j_local2), ldda);
#endif
#if defined (ENABLE_TIMER)
magma_device_sync();
tmnp = magma_wtime() - tmnp;
ttot_mnp += tmnp;
#endif
}
// send next panel to cpu
magma_queue_sync( queues[d][STREAM_ID(j_local2)] ); // make sure lookahead is done
tmppanel = ((blkid+1) % 2 == 0) ? tmppanel0 : tmppanel1;
magma_dgetmatrix_async(rows-nb, nb,
dlA(d, j+nb, j_local2), ldda,
tmppanel(j+nb), ldpanel,
queues[d][0] );
tmppanel = (blkid % 2 == 0) ? tmppanel0 : tmppanel1;
offset = j + nb + nb;
row_offset = nb;
} else {
offset = j + nb;
row_offset = 0;
}
if (n - offset > 0) {
// syrk on multiple gpu
for (d = 0; d < ngpu; d++ ) {
if ( d == id ) {
dlpanels[d] = dlA(d, j + nb + row_offset, j_local);
} else {
dlpanels[d] = dlP(d, row_offset, 0);
}
}
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
therk[d] = magma_wtime();
}
#endif
//magmablasSetKernelStream( queues[d] );
//magma_dsyrk( MagmaLower, MagmaNoTrans, n - offset, nb,
// d_neg_one, dlpanel, ldda,
// d_one, &d_lA[d][offset + offset*ldda], ldda );
#ifdef DSYRK_ON_DIAG
magma_dsyrk_mgpu
#else
magma_dsyrk_mgpu2
#endif
(ngpu, MagmaLower, MagmaNoTrans,
nb, n - offset, nb,
d_neg_one, dlpanels, ldda, 0,
d_one, d_lA, ldda, offset,
nqueue, queues );
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_device_sync();
therk[d] = magma_wtime() - therk[d];
ttot_herk[d] += therk[d];
}
#endif
}
prevtrsmrows = trsmrows;
#if defined (ENABLE_TIMER)
ttot += (tcnp+tcchol+tctrsm+therk[0]+therk[1]+therk[2]+tctm+tmnp);
printf("%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(%d) %10.3lf\n",
j, nb, rows, tmtc,
tcnp, // gemm
tcchol, // potrf
tctrsm, // trsm
(tcchol + tctrsm),
(tmtc+tcnp+tcchol+tctrsm),
therk[0], therk[1], therk[2], therk[3], // syrk
tctm, // copy panel to GPU
tmnp, // lookahead on GPU
(id + 1) % ngpu,
(tcnp+tcchol+tctrsm+therk[0]+therk[1]+therk[2]+tctm+tmnp));
fflush(0);
#endif
}
}
for( d = 0; d < ngpu; d++ ) {
magma_setdevice(d);
for( id=0; id < nqueue; id++ ) {
magma_queue_sync( queues[d][id] );
}
}
#if defined (ENABLE_TIMER)
printf("\n%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(-) %10.3lf\n",
n, n, 0, ttot_mtc,
ttot_cnp, // gemm
ttot_cchol, // potrf
ttot_ctrsm, // trsm
(ttot_cchol + ttot_ctrsm),
(ttot_mtc+ttot_cnp+ttot_cchol+ttot_ctrsm),
ttot_herk[0], ttot_herk[1], ttot_herk[2], ttot_herk[3], // syrk
ttot_ctm, // copy panel to GPU
ttot_mnp, // lookahead on GPU
(ttot_cnp+ttot_cchol+ttot_ctrsm+ttot_herk[0]+ttot_herk[1]+ttot_herk[2]+ttot_ctm+ttot_mnp));
printf("%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(-) %10.3lf (ratio)\n",
n, n, 0, ttot_mtc/ttot,
ttot_cnp/ttot, // gemm
ttot_cchol/ttot, // potrf
ttot_ctrsm/ttot, // trsm
(ttot_cchol + ttot_ctrsm)/ttot,
(ttot_mtc+ttot_cnp+ttot_cchol+ttot_ctrsm)/ttot,
ttot_herk[0]/ttot, ttot_herk[1]/ttot, ttot_herk[2]/ttot, ttot_herk[3]/ttot, // syrk
ttot_ctm/ttot, // copy panel to GPU
ttot_mnp/ttot, // lookahead on GPU
(ttot_cnp+ttot_cchol+ttot_ctrsm+ttot_herk[0]+ttot_herk[1]+ttot_herk[2]+ttot_ctm+ttot_mnp)/ttot);
#endif
// cholesky for the last block
if (j < n && *info == 0) {
rows = n - j;
id = (j / nb) % ngpu;
// Set the local index where the current panel is
j_local = j / (nb * ngpu) * nb;
magma_setdevice(id);
#if defined (ENABLE_TIMER)
tset = magma_wtime();
#endif
magma_dgetmatrix(rows, rows, dlA(id, j, j_local), ldda, panel(j), ldpanel);
lapackf77_dpotrf(MagmaLowerStr, &rows, panel(j), &ldpanel, info);
magma_dsetmatrix(rows, rows, panel(j), ldpanel, dlA(id, j, j_local), ldda);
#if defined (ENABLE_TIMER)
tset = magma_wtime() - tset;
#endif
}
#if defined (ENABLE_TIMER)
printf( " matrix_get,set: %10.3lf %10.3lf -> %10.3lf\n",tget,tset,ttot+tget+tset );
#endif
} // end of else not upper
// clean up
for( d = 0; d < ngpu; d++ ) {
magma_setdevice(d);
for( j=0; j < nqueue; j++ ) {
magma_queue_destroy( queues[d][j] );
}
magma_free( d_lP[d] );
}
} // end of not lapack
// free workspace
magma_free_pinned( panel );
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_dpotrf_mgpu_right */
