extern "C" magma_int_t
magma_cpotrf2_msub(
magma_int_t num_subs, 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,
magmaFloatComplex_ptr *d_lA, size_t d_lA_offset, magma_int_t ldda,
magmaFloatComplex_ptr *d_lP, magma_int_t lddp,
magmaFloatComplex *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
=======
CPOTRF computes the Cholesky factorization of a complex Hermitian
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) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the Hermitian 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.
===================================================================== */
int tot_subs = num_subs*num_gpus;
magma_int_t j, jb, nb0, nb2, dd, d, id, j_local, j_local2;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.0;
float d_neg_one = -1.0;
magmaFloatComplex_ptr dlpanel;
size_t dlpanel_offset;
magma_int_t n_local[MagmaMaxSubs * MagmaMaxGPUs], ldpanel;
// initialize trace
trace_init(1, num_gpus, 2, queues);
*info = 0;
if ( (uplo != MagmaUpper) && (uplo != MagmaLower) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if ((uplo != MagmaUpper) && tot_subs*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<tot_subs; d++) {
/* local-n and local-ld */
if (uplo == MagmaUpper) {
n_local[d] = ((n/nb)/tot_subs)*nb;
if (d < (n/nb)%tot_subs)
n_local[d] += nb;
else if (d == (n/nb)%tot_subs)
n_local[d] += n%nb;
} else {
n_local[d] = ((m/nb)/tot_subs)*nb;
if (d < (m/nb)%tot_subs)
n_local[d] += nb;
else if (d == (m/nb)%tot_subs)
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)%tot_subs;
/* Set the local index where the current panel is */
j_local = j/(nb*tot_subs);
jb = min(nb, (m-j));
if (j > 0) {
// Wait for the column on CPU
magma_queue_sync(queues[2*(id%num_gpus)]);
/* broadcast off-diagonal column to all gpus */
d = (j/nb+1)%num_gpus;
for (dd=0; dd<num_gpus; dd++) {
if (d != id%num_gpus) {
magma_csetmatrix_async( j, jb,
Aup(0,j), lda,
dlP(d,jb,0,id%num_gpus), lddp,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-col") );
}
d = (d+1)%num_gpus;
}
/* Update the current diagonal block */
trace_gpu_start(id%num_gpus, 1, "herk", "herk");
magma_cherk(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%num_gpus)+1]);
magma_queue_sync(queues[2*(id%num_gpus)+1]); // Wait for syrk
}
/* Send the diagonal to cpu */
magma_cgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "get", "get-diag") );
if (j > 0) {
/* Compute the local block column of the panel. */
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
j_local2 = j_local+1;
if (d > id) j_local2 --;
nb0 = nb*j_local2;
if (n_local[d] > nb0) {
if (d%num_gpus != id%num_gpus) {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlP_offset(jb, 0, id%num_gpus);
ldpanel = lddp;
/* Wait for the offdiagonal column */
if (dd < num_gpus) magma_queue_sync(queues[2*(d%num_gpus)]);
} else {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(0, nb*j_local);
ldpanel = ldda;
}
/* update the panel */
trace_gpu_start(d%num_gpus, 1, "gemm", "gemm");
magma_cgemm(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%num_gpus)+1]);
}
d = (d+1)%tot_subs;
}
}
/* factor the diagonal */
magma_queue_sync( queues[2*(id%num_gpus)] ); // wait for the diagonal
trace_cpu_start(0, "potrf", "potrf");
lapackf77_cpotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
trace_cpu_end(0);
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%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = d_lP[d];
dlpanel_offset = dlP_offset(0, 0, id%num_gpus);
ldpanel = lddp;
}
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, dlpanel_offset, ldpanel,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-diag"));
d = (d+1)%num_gpus;
}
} else {
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "set", "set-diag") );
}
/* panel-factorize the off-diagonal */
if ((j+jb) < n) {
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
/* next column */
j_local2 = j_local+1;
if (d > id) j_local2--;
if (d%num_gpus == id%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlP_offset(0, 0, id%num_gpus);
ldpanel = lddp;
}
nb2 = n_local[d]-nb*j_local2;
nb0 = min(nb, nb2);
if (dd < num_gpus) magma_queue_sync( queues[2*(d%num_gpus)] ); // wait for the diagonal
if (j+jb < m && d == (j/nb+1)%tot_subs) {
/* owns the next column, look-ahead the column */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2), ldda,
queues[2*(d%num_gpus)+1] );
/* send the column to cpu */
magma_queue_sync(queues[2*(d%num_gpus)+1]); // wait for lookahead
magma_cgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
queues[2*(d%num_gpus)],
trace_gpu_event(d%num_gpus, 0, "get", "get-col") );
/* update the remaining blocks */
nb2 = nb2 - nb0;
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
queues[2*(d%num_gpus)+1] );
} else if (nb2 > 0) {
/* update the entire trailing matrix */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2), ldda,
queues[2*(d%num_gpus)+1] );
}
d = (d+1)%tot_subs;
}
}
}
} 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)%tot_subs;
/* Set the local index where the current panel is */
j_local = j/(nb*tot_subs);
jb = min(nb, (n-j));
if (j > 0) {
if (num_gpus > 1) {
// Wait for the row on CPU to broadcast
magma_queue_sync(queues[2*(id%num_gpus)]);
}
/* broadcast off-diagonal row to all the GPUs */
d = (j/nb+1)%num_gpus;
for (dd=0; dd<num_gpus; dd++) {
if (d != id%num_gpus) {
/* send it to GPU-d */
magma_csetmatrix_async( jb, j,
Alo(j,0), lda,
dlPT(d,0,jb,id%num_gpus), nb,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-row") );
}
d = (d+1)%num_gpus;
}
/* Update the current diagonal block */
trace_gpu_start(id%num_gpus, 1, "herk", "herk");
magma_cherk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda,
queues[2*(id%num_gpus)+1]);
magma_queue_sync(queues[2*(id%num_gpus)+1]); // wait for syrk
}
/* send the diagonal to cpu */
magma_cgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "get", "get") );
/* update the offdiagonal blocks */
if (j > 0) {
/* compute the block-rows of the panel */
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
j_local2 = j_local+1;
if (d > id) j_local2 --;
nb0 = nb*j_local2;
if (nb0 < n_local[d]) {
if (d%num_gpus != id%num_gpus) {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlPT_offset(0, jb, id%num_gpus);
ldpanel = nb;
/* Wait for offdiagonal row */
if (dd < num_gpus) magma_queue_sync(queues[2*(d%num_gpus)]);
} else {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(nb*j_local, 0);
ldpanel = ldda;
}
/* Update the panel */
trace_gpu_start(d%num_gpus, 1, "gemm", "gemm");
magma_cgemm( 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[2*(d%num_gpus)+1]);
}
d = (d+1)%tot_subs;
}
}
/* factor the diagonal */
magma_queue_sync( queues[2*(id%num_gpus)] );
trace_cpu_start(0, "potrf", "potrf");
lapackf77_cpotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
trace_cpu_end(0);
if (*info != 0) {
printf( " cpotrf returned %d (id=%d,j=%d,j_local=%d,jb=%d)\n",*info,id,j,j_local,jb );
*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%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = d_lP[d];
dlpanel_offset = dlPT_offset(0, 0, id%num_gpus);
ldpanel = nb;
}
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, dlpanel_offset, ldpanel,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-diag") );
d = (d+1)%num_gpus;
}
} else {
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "set", "set-diag") );
}
/* factorize off-diagonal blocks */
if ((j+jb) < m) {
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
/* next column */
j_local2 = j_local+1;
if (d > id) j_local2--;
if (d%num_gpus == id%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlPT_offset(0, 0, id%num_gpus);
ldpanel = nb;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2 );
// wait for the diagonal
if (dd < num_gpus) magma_queue_sync(queues[2*(d%num_gpus)]);
if (j+jb < n && d == (j/nb+1)%tot_subs) {
/* owns the next column, look-ahead the column */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2, j), ldda,
queues[2*(d%num_gpus)+1]);
/* send the column to cpu */
magma_queue_sync( queues[2*(d%num_gpus)+1] ); // wait for lookahead
magma_cgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
queues[2*(d%num_gpus)],
trace_gpu_event(d%num_gpus, 0, "get", "get") );
/* update the remaining blocks */
nb2 = nb2 - nb0;
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
queues[2*(d%num_gpus)+1]);
} else if (nb2 > 0) {
/* update the entire trailing matrix */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2, j), ldda,
queues[2*(d%num_gpus)+1]);
}
d = (d+1)%tot_subs;
}
}
}
} /* end of else not upper */
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_queue_sync( queues[2*d] );
magma_queue_sync( queues[2*d+1] );
}
trace_finalize("cpotrf_msub.svg", "trace.css");
return *info;
} /* magma_cpotrf2_msub */
/**
Purpose
-------
CLAUUM 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 = MagmaUpper then the upper triangle of the result is stored,
overwriting the factor U in dA.
If UPLO = MagmaLower 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
---------
@param[in]
uplo magma_uplo_t
Specifies whether the triangular factor stored in the array dA
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]
dA REAL array on the GPU, dimension (LDDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = MagmaUpper, the upper triangle of dA is
overwritten with the upper triangle of the product U * U';
if UPLO = MagmaLower, the lower triangle of dA is overwritten with
the lower triangle of the product L' * L.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -k, the k-th argument had an illegal value
@ingroup magma_cposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_clauum_gpu(magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda, magma_int_t *info)
{
#define dA(i, j) (dA + (j)*ldda + (i))
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t nb, i, ib;
float d_one = MAGMA_D_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *work;
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_cpotrf_nb(n);
if (MAGMA_SUCCESS != magma_cmalloc_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_cgetmatrix( n, n, dA, ldda, work, n );
lapackf77_clauum(uplo_, &n, work, &n, info);
magma_csetmatrix( 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_ctrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
magma_cgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_clauum(MagmaUpperStr, &ib, work, &ib, info);
magma_csetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if (i+ib < n) {
magma_cgemm( 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);
magma_cherk( 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 += nb) {
ib=min(nb,(n-i));
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
magma_cgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_clauum(MagmaLowerStr, &ib, work, &ib, info);
magma_csetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if (i+ib < n) {
magma_cgemm( 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);
magma_cherk( MagmaLower, MagmaConjTrans, 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;
}
extern "C" void
magma_cherk_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
float alpha,
magmaFloatComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
float beta,
magmaFloatComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t nqueue, magma_queue_t queues[][10])
{
#define dB(id, i, j) (dB[(id)]+(j)*lddb + (i)+b_offset)
#define dC(id, i, j) (dC[(id)]+(j)*lddc + (i))
#define STREAM_ID(i) (nqueue > 1 ? 1+((i)/nb)%(nqueue-1) : 0)
magma_int_t i, id, ib, ii, kk, n1;
magmaFloatComplex z_alpha = MAGMA_C_MAKE(alpha,0.0);
magmaFloatComplex z_beta = MAGMA_C_MAKE(beta, 0.0);
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
/* cher2k on diagonal block */
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_cherk(uplo, trans, ib, k,
alpha, dB(id, i, 0 ), lddb,
beta, dC(id, i+c_offset, ii), lddc);
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
z_alpha, dB(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
z_beta, dC(id, 0, ii), lddc);
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
/* cgemm on off-diagonal blocks */
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
z_alpha, dB(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
z_beta, dC(id, i+c_offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
// TODO why not sync?
//for( id=0; id < ngpu; id++ ) {
// magma_setdevice(id);
// //for( kk=0; kk < nqueue; kk++ )
// // magma_queue_sync( queues[id][kk] );
//}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
extern "C" magma_int_t
magma_cpotrf3_mgpu(magma_int_t 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,
magmaFloatComplex *d_lA[], magma_int_t ldda,
magmaFloatComplex *d_lP[], magma_int_t lddp,
magmaFloatComplex *a, magma_int_t lda, magma_int_t h,
magma_queue_t stream[][3], magma_event_t event[][5],
magma_int_t *info )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
Auxiliary subroutine for cpotrf2_ooc. It is multiple gpu interface to compute
Cholesky of a "rectangular" matrix.
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) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the Hermitian 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
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};
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.0;
float d_neg_one = -1.0;
int upper = lapackf77_lsame(uplo_, "U");
magmaFloatComplex *dlpanel;
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
const magma_int_t stream1 = 0, stream2 = 1, stream3 = 2;
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
/* used by ctrsm_work */
int trsm_nb = 128;
int trsm_n = trsm_nb*((nb+trsm_nb-1)/trsm_nb);
magmaFloatComplex *d_dinvA[MagmaMaxGPUs];
magmaFloatComplex *d_x[MagmaMaxGPUs];
#define dinvA(d,j) &(d_dinvA[(d)][(j)*trsm_nb*trsm_n])
#define dx(d,j) &(d_x[(d)][(j)*nb*m])
/*
* Allocate device memory for the inversed diagonal blocks, size=N*BLOCK_SIZE
*/
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
if ( (MAGMA_SUCCESS != magma_cmalloc( &d_dinvA[d], 2*trsm_nb*trsm_n )) ||
(MAGMA_SUCCESS != magma_cmalloc( &d_x[d], 2*nb*(upper ? n : m) )) ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
magma_setdevice(0);
#endif
*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;
}
}
/* == initialize the trace */
trace_init( 1, num_gpus, 3, (CUstream_st**)stream );
if (upper)
{
/* ---------------------------------------------- */
/* 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; // right now, we have num_gpu buffers, so id and buf are the same..
/* 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_cherk(MagmaUpper, MagmaConjTrans, 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_cgetmatrix_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; // number of local columns in the panel, while jb is panel-size (number of rows)
if( n_local[d] > nb0 ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
if( d == id ) {
dlpanel = dlA(d,0,nb*j_local);
ldpanel = ldda;
// the GPU owns the row from start, and no need of synch.
//magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
} else {
dlpanel = dlP(d,nb,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_cgemm(MagmaConjTrans, 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_cpotrf(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_csetmatrix_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_csetmatrix_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;
magma_setdevice(d);
if( j+jb < m && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead next block on stream1 */
nb0 = min(nb, nb2);
magmablasSetKernelStream(stream[d][stream1]);
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for gemm update
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb, trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb0,jb, dx(d,0),nb0 );
magmablas_ctrsm_work( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaLeft, MagmaUpper,
MagmaConjTrans, 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(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,0),nb2 );
magmablas_ctrsm_work( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaLeft, MagmaUpper,
MagmaConjTrans, 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 */
/* ========================================================== */
if( j+jb < m ) {
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. */
int d2, buf2;
magma_setdevice(d);
/* lookahead done */
magma_queue_wait_event( stream[d][stream3], event[d][4] );
trace_gpu_start( d, stream3, "comm", "row to CPU" );
magma_cgetmatrix_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] );
/* 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_csetmatrix_async( j+jb, nb0,
Aup(0,j+jb), lda,
dlP(d2,nb,0,buf2), lddp,
stream[d2][stream3] );
trace_gpu_end( d2, stream3 );
magma_event_record( event[d2][0], stream[d2][stream3] );
}
}
/* =========================== */
/* update the remaining blocks */
nb2 = n_local[d]-(nb*j_local2 + nb0);
if( nb2 > 0 ) {
if( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
trace_gpu_start( d, stream2, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
int flag = 0;
if (flag == 0) {
magma_queue_wait_event( stream[d][stream2], event[d][4] ); // lookahead -> diagonal inversion
} else {
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,flag),trsm_nb );
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
}
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,1),nb2 );
magmablas_ctrsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
flag, dinvA(d,flag), dx(d,1) );
#else
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for cholesky factor
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda);
#endif
trace_gpu_end( d, stream2 );
}
}
} /* end of ctrsm */
} /* end of for j=1, .., n */
} 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));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
magma_cherk(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_cgetmatrix_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,nb,buf);
ldpanel = nb;
magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
}
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
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_cpotrf(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_csetmatrix_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_csetmatrix_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 */
if ( j > 0 ) 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(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb, trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb0,jb, dx(d,0),nb0 );
magmablas_ctrsm_work( MagmaRight, MagmaLower,
MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaRight, MagmaLower,
MagmaConjTrans, 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(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,0),nb2 );
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
}
d = (d+1)%num_gpus;
} /* end for d */
/* ========================================================== */
if( j+jb < n ) {
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. */
int d2, buf2;
//#define CPOTRF_DEVICE_TO_DEVICE
#ifdef CPOTRF_DEVICE_TO_DEVICE
// lookahead done
/* broadcast the rows to gpus */
buf2 = ((j+jb)/nb)%num_gpus;
for( d2=0; d2<num_gpus; d2++ ) {
magma_setdevice(d2);
magma_queue_wait_event( stream[d2][stream3], event[d][4] );
if( d2 != d ) {
magma_ccopymatrix_async( nb0, j+jb,
dlPT(d2,0,nb,buf2), nb, // first nbxnb reserved for diagonal block
dlA(d, nb*j_local2, 0), ldda,
stream[d2][stream3] );
magma_event_record( event[d2][0], stream[d2][stream3] );
} else {
magma_cgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
stream[d][stream3] );
}
}
#else
// lookahead done
magma_setdevice(d);
magma_queue_wait_event( stream[d][stream3], event[d][4] );
magma_cgetmatrix_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_csetmatrix_async( nb0, j+jb,
Alo(j+jb,0), lda,
dlPT(d2,0,nb,buf2), nb, // first nbxnb reserved for diagonal block
stream[d2][stream3] );
magma_event_record( event[d2][0], stream[d2][stream3] );
}
}
#endif
/* =================================== */
/* updates remaining blocks on stream2 */
nb2 = n_local[d] - (j_local2*nb + nb0);
if( nb2 > 0 ) {
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d,0,0,buf);
ldpanel = nb;
}
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
/* update the remaining blocks in the column */
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
int flag = 0;
if (flag == 0) {
magma_queue_wait_event( stream[d][stream2], event[d][4] ); // lookahead -> diagonal inversion
} else {
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,flag),trsm_nb );
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
}
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,1),nb2 );
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
flag, dinvA(d,flag), dx(d,1) );
#else
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, 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( "cpotrf.svg","trace.css" );
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j<3; j++ ) {
magma_queue_sync( stream[d][j] );
}
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magma_free( d_dinvA[d] );
magma_free( d_x[d] );
#endif
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
return *info;
} /* magma_cpotrf_mgpu */
extern "C" magma_int_t
magma_cpotrf_gpu(char uplo, magma_int_t n,
magmaFloatComplex *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
=======
CPOTRF computes the Cholesky factorization of a complex Hermitian
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.
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) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the Hermitian 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
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};
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *work;
float d_one = 1.0;
float 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_cpotrf_nb(n);
if (MAGMA_SUCCESS != magma_cmalloc_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_cgetmatrix_async( n, n, dA, ldda, work, n, stream[1] );
magma_queue_sync( stream[1] );
lapackf77_cpotrf(uplo_, &n, work, &n, info);
magma_csetmatrix_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_cherk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[0] );
if ( (j+jb) < n) {
/* Compute the current block row. */
magma_cgemm(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_cpotrf(MagmaUpperStr, &jb, work, &jb, info);
magma_csetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[1] );
if (*info != 0) {
*info = *info + j;
break;
}
if ( (j+jb) < n) {
magma_ctrsm( 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_cherk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda);
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[0] );
if ( (j+jb) < n) {
magma_cgemm( 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_cpotrf(MagmaLowerStr, &jb, work, &jb, info);
magma_csetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[1] );
if (*info != 0) {
*info = *info + j;
break;
}
if ( (j+jb) < n) {
magma_ctrsm(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 (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
return *info;
} /* magma_cpotrf_gpu */
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaFloatComplex *A, *B, *C, *C2, *LU;
magmaFloatComplex_ptr dA, dB, dC1, dC2;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.5, 0.1 );
magmaFloatComplex beta = MAGMA_C_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float 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_int_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 itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
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 = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_cmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_cmalloc( &dA, size ); assert( err == 0 );
err = magma_cmalloc( &dB, size ); assert( err == 0 );
err = magma_cmalloc( &dC1, size ); assert( err == 0 );
err = magma_cmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test CSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_csetmatrix( m, n, A, ld, dA, ld );
magma_csetmatrix( m, n, A, ld, dB, ld );
magma_cswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_cswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_cgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_clange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "cswap diff %.2g\n", error );
}
else {
printf( "cswap skipped for n < 3\n" );
}
// ----- test ICAMAX
// get argmax of column of A
magma_csetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_icamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIcamax( opts.handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "icamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test CGEMV
// 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_csetmatrix( m, n, A, ld, dA, ld );
magma_csetvector( maxn, B, 1, dB, 1 );
magma_csetvector( maxn, C, 1, dC1, 1 );
magma_csetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_cgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCgemv( opts.handle, cublas_trans_const(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] == MagmaNoTrans ? m : n);
cublasCaxpy( opts.handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMV( m, n ) / 1e9;
printf( "cgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CHEMV
// 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_csetmatrix( m, m, A, ld, dA, ld );
magma_csetvector( m, B, 1, dB, 1 );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_chemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasChemv( opts.handle, cublas_uplo_const(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
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMV( m ) / 1e9;
printf( "chemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CTRSV
// 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_clacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_cgetrf( &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_csetmatrix( m, m, LU, ld, dA, ld );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ctrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCtrsv( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ctrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test CGEMM
// 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] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_csetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_csetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cgemm( 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 );
cublasCgemm( opts.handle, cublas_trans_const(trans[ia]), cublas_trans_const(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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMM( m, n, k ) / 1e9;
printf( "cgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHEMM
// 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_csetmatrix( m, m, A, ld, dA, ld );
magma_csetmatrix( m, n, B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_chemm( 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 );
cublasChemm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMM( side[is], m, n ) / 1e9;
printf( "chemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHERK
// 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_csetmatrix( n, k, A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCherk( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHERK( k, n ) / 1e9;
printf( "cherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHER2K
// 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] == MagmaNoTrans);
magma_csetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cher2k( 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 );
cublasCher2k( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHER2K( k, n ) / 1e9;
printf( "cher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CTRMM
// 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] == MagmaLeft);
magma_csetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasCtrmm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRMM( side[is], m, n ) / 1e9;
printf( "ctrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test CTRSM
// 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] == MagmaLeft);
magma_csetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrsm( 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 );
cublasCtrsm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(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
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( side[is], m, n ) / 1e9;
printf( "ctrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(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 );
fflush( stdout );
}
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();
int status = (total_error != 0.);
return status;
}
extern "C" magma_int_t
magma_cpotrf_m(magma_int_t num_gpus0, char uplo, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CPOTRF_OOC computes the Cholesky factorization of a complex Hermitian
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 not fit entirely in the GPU memory.
The factorization has the form
A = U**H * U, if UPLO = 'U', or
A = 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 A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX 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**H * U or A = L * L**H.
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.
===================================================================== */
/* Local variables */
float d_one = 1.0;
float d_neg_one = -1.0;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
char uplo_[2] = {uplo, 0};
int upper = lapackf77_lsame(uplo_, "U");
magmaFloatComplex *dwork[MagmaMaxGPUs], *dt[MagmaMaxGPUs];
magma_int_t ldda, lddla, nb, iinfo, n_local[MagmaMaxGPUs], J2, d, num_gpus;
magma_int_t j, jj, jb, J, JB, NB, MB, h;
magma_queue_t stream[MagmaMaxGPUs][3];
magma_event_t event[MagmaMaxGPUs][5];
#ifdef ROW_MAJOR_PROFILE
magma_timestr_t start, end, start0, end0;
float chol_time = 1.0;
#endif
*info = 0;
if ((! upper) && (! lapackf77_lsame(uplo_, "L"))) {
*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( num_gpus0 > n/nb ) {
num_gpus = n/nb;
if( n%nb != 0 ) num_gpus ++;
} else {
num_gpus = num_gpus0;
}
//ldda = ((n+31)/32)*32;
ldda = ((n+nb-1)/nb)*nb;
lddla = ((nb*((n+nb*num_gpus-1)/(nb*num_gpus))+31)/32)*32;
/* figure out NB */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
MB = n; /* number of rows in the big panel */
NB = (magma_int_t)((0.8*freeMem-max(2,num_gpus)*nb*ldda-(n+nb)*nb)/lddla); /* number of columns in the big panel */
//NB = min(5*nb,n);
if( NB >= n ) {
#ifdef CHECK_CPOTRF_OOC
printf( " * still fit in GPU memory.\n" );
#endif
NB = n;
} else {
#ifdef CHECK_CPOTRF_OOC
printf( " * don't fit in GPU memory.\n" );
#endif
NB = (NB/nb) * nb; /* making sure it's devisable by nb */
}
#ifdef CHECK_CPOTRF_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,(float)freeMem );
else printf( " * running in in-core mode (n=%d, NB=%d, nb=%d, lddla=%d, freeMem=%.2e).\n",n,NB,nb,lddla,(float)freeMem );
fflush(stdout);
#endif
for (d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_cmalloc( &dt[d], NB*lddla + max(2,num_gpus)*nb*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork[d] = &dt[d][max(2,num_gpus)*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);
#ifdef ROW_MAJOR_PROFILE
start0 = get_current_time();
#endif
if (nb <= 1 || nb >= n) {
lapackf77_cpotrf(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( num_gpus0 > (n-J)/nb ) {
num_gpus = (n-J)/nb;
if( (n-J)%nb != 0 ) num_gpus ++;
} else {
num_gpus = num_gpus0;
}
/* load the new big-panel by block-rows */
magma_chtodpo( num_gpus, &uplo, JB, n, J, J, nb, a, lda, dwork, NB, stream, &iinfo);
#ifdef ROW_MAJOR_PROFILE
start = get_current_time();
#endif
/* update with the previous big-panels */
for( j=0; j<J; j+=nb ) {
/* upload the diagonal of the block column (broadcast to all GPUs) */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magma_csetmatrix_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)%num_gpus;
magma_setdevice(d);
jb = min(nb, n-jj);
magma_csetmatrix_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<num_gpus; 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)%num_gpus;
J2 = jj/(nb*num_gpus);
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_cgemm( 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_cherk(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<num_gpus; d++ ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][2]);
/* local number of columns in the big panel */
n_local[d] = ((n-J)/(nb*num_gpus))*nb;
if (d < ((n-J)/nb)%num_gpus)
n_local[d] += nb;
else if (d == ((n-J)/nb)%num_gpus)
n_local[d] += (n-J)%nb;
/* subtracting the local number of columns in the diagonal */
J2 = nb*(JB/(nb*num_gpus));
if( d < (JB/nb)%num_gpus ) J2+=nb;
n_local[d] -= J2;
magma_cgemm( 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<num_gpus; 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 two streams
//magma_cpotrf2_mgpu(num_gpus, uplo, JB, n-J, J, J, nb,
// dwork, NB, dt, ldda, a, lda, h, stream, event, &iinfo);
// using three streams
magma_cpotrf3_mgpu(num_gpus, uplo, JB, n-J, J, J, nb,
dwork, NB, dt, ldda, a, lda, h, stream, event, &iinfo);
if( iinfo != 0 ) {
*info = J+iinfo;
break;
}
#ifdef ROW_MAJOR_PROFILE
end = get_current_time();
chol_time += GetTimerValue(start, end);
#endif
/* upload the off-diagonal (and diagonal!!!) big panel */
magma_cdtohpo(num_gpus, &uplo, JB, n, J, J, nb, NB, a, lda, dwork, NB, stream, &iinfo);
//magma_cdtohpo(num_gpus, &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( num_gpus0 > (n-J)/nb ) {
num_gpus = (n-J)/nb;
if( (n-J)%nb != 0 ) num_gpus ++;
} else {
num_gpus = num_gpus0;
}
/* load the new big-panel by block-columns */
magma_chtodpo( num_gpus, &uplo, n, JB, J, J, nb, a, lda, dwork, lddla, stream, &iinfo);
/* update with the previous big-panels */
#ifdef ROW_MAJOR_PROFILE
start = get_current_time();
#endif
for( j=0; j<J; j+=nb ) {
/* upload the diagonal of big panel */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magma_csetmatrix_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)%num_gpus;
magma_setdevice(d);
jb = min(nb, n-jj);
magma_csetmatrix_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<num_gpus; 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)%num_gpus;
J2 = jj/(nb*num_gpus);
magma_setdevice(d);
magmablasSetKernelStream(stream[d][J2%2]);
J2 = nb*J2;
jb = min(nb,JB-jj);
magma_cgemm( 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_cherk(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<num_gpus; d++ ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][2]);
/* local number of columns in the big panel */
n_local[d] = (((n-J)/nb)/num_gpus)*nb;
if (d < ((n-J)/nb)%num_gpus)
n_local[d] += nb;
else if (d == ((n-J)/nb)%num_gpus)
n_local[d] += (n-J)%nb;
/* subtracting local number of columns in diagonal */
J2 = nb*(JB/(nb*num_gpus));
if( d < (JB/nb)%num_gpus ) J2+=nb;
n_local[d] -= J2;
magma_cgemm( 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<num_gpus; 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 two streams
//magma_cpotrf2_mgpu(num_gpus, uplo, n-J, JB, J, J, nb,
// dwork, lddla, dt, ldda, a, lda, h, stream, event, &iinfo);
// using three streams
magma_cpotrf3_mgpu(num_gpus, uplo, n-J, JB, J, J, nb,
dwork, lddla, dt, ldda, a, lda, h, stream, event, &iinfo);
if( iinfo != 0 ) {
*info = J+iinfo;
break;
}
#ifdef ROW_MAJOR_PROFILE
end = get_current_time();
chol_time += GetTimerValue(start, end);
#endif
/* upload the off-diagonal big panel */
magma_cdtohpo( num_gpus, &uplo, n, JB, J, J, nb, JB, a, lda, dwork, lddla, stream, &iinfo);
} /* end of for J */
} /* if upper */
} /* if nb */
#ifdef ROW_MAJOR_PROFILE
end0 = get_current_time();
#endif
if( num_gpus0 > n/nb ) {
num_gpus = n/nb;
if( n%nb != 0 ) num_gpus ++;
} else {
num_gpus = num_gpus0;
}
for (d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j<3; j++ ) {
if( stream[d][j] != NULL ) magma_queue_destroy( stream[d][j] );
}
magma_free( dt[d] );
for( j=0; j<5; j++ ) {
magma_event_destroy( event[d][j] );
}
}
magma_setdevice(0);
#ifdef ROW_MAJOR_PROFILE
printf("\n n=%d NB=%d nb=%d\n",n,NB,nb);
printf(" Without memory allocation: %f / %f = %f GFlop/s\n",
FLOPS_CPOTRF(n)/1000000, GetTimerValue(start0, end0),
FLOPS_CPOTRF(n)/(1000000*GetTimerValue(start0, end0)));
printf(" Performance %f / %f = %f GFlop/s\n",
FLOPS_CPOTRF(n)/1000000, chol_time,
FLOPS_CPOTRF(n)/(1000000*chol_time));
#endif
return *info;
} /* magma_cpotrf_ooc */
/**
Purpose
-------
CPOTRF computes the Cholesky factorization of a complex Hermitian
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]
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 COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the Hermitian 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_cposv_comp
********************************************************************/
extern "C" magma_int_t
magma_cpotrf2_mgpu(int 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,
magmaFloatComplex **d_lA, magma_int_t ldda,
magmaFloatComplex **d_lP, magma_int_t lddp,
magmaFloatComplex *A, magma_int_t lda, magma_int_t h,
magma_queue_t stream[][3], magma_event_t event[][5],
magma_int_t *info )
{
#define Alo(i, j) (A + ((j)+off_j)*lda + (nb*(((i)/nb)%h)+off_i))
#define Aup(i, j) (A + (nb*(((j)/nb)%h)+off_j)*lda + (i+off_i))
#define dlA(id, i, j) (d_lA[(id)] + (j)*ldda + (i))
#define dlP(id, i, j, k) (d_lP[(id)] + (k)*nb*lddp + (j)*lddp + (i))
#define dlPT(id, i, j, k) (d_lP[(id)] + (k)*nb*lddp + (j)*nb + (i))
magma_int_t j, jb, nb0, nb2, dd, d, id, j_local, j_local2, buf;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.0;
float d_neg_one = -1.0;
int upper = (uplo == MagmaUpper);
magmaFloatComplex *dlpanel;
//magma_event_t event0[MagmaMaxGPUs], // syrk
// event1[MagmaMaxGPUs], // send off-diagonal
// event2[MagmaMaxGPUs], // send diagonal
// event3[MagmaMaxGPUs]; // trsm
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
int stream0 = 0, stream1 = 1;
#ifdef CTRSM_WORK
magmaFloatComplex *d_dinvA[MagmaMaxGPUs][2], *d_x[MagmaMaxGPUs][2]; /* used by ctrsm_work */
#endif
*info = 0;
if (! upper && uplo != MagmaLower) {
*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;
}
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_setdevice(0);
/* == initialize the trace */
trace_init( 1, num_gpus, 3, (magma_queue_t*)stream );
/* Use blocked code. */
if (upper) {
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
#if defined(PRECISION_d) && defined(CTRSM_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++ ) {
magma_cmalloc( &d_dinvA[d][j], nb*nb );
magma_cmalloc( &d_x[d][j], n*nb );
cudaMemset(d_dinvA[d][j], 0, nb*nb*sizeof(magmaFloatComplex));
cudaMemset(d_x[d][j], 0, n*nb*sizeof(magmaFloatComplex));
}
}
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));
if ( j > 0 ) {
/* needed on pluto... */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] ); // 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_setdevice(d);
/* wait for it on CPU */
magma_queue_wait_event( stream[d][stream0], event[id][1] );
/* send it to GPU */
trace_gpu_start( d, stream0, "comm", "rows to GPUs" );
magma_csetmatrix_async( j, jb,
Aup(0,j), lda,
dlP(d,jb,0,buf), lddp,
stream[d][stream0] );
trace_gpu_end( d, stream0 );
magma_event_record( event[d][1], stream[d][stream0] );
}
d = (d+1)%num_gpus;
}
}
/* Update the current diagonal block */
magma_setdevice(id);
if ( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
trace_gpu_start( id, stream1, "syrk", "syrk" );
magma_cherk(MagmaUpper, MagmaConjTrans, 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 );
magma_event_record( event[id][0], stream[id][stream1] );
}
/* send the diagonal to cpu */
magma_queue_wait_event( stream[id][stream0], event[id][0] ); // wait for syrk
trace_gpu_start( id, stream0, "comm", "D to CPU" );
magma_cgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
stream[id][stream0] );
trace_gpu_end( id, stream0 );
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 ) {
//magma_queue_sync( stream[id][3] );
dlpanel = dlP(d, jb, 0, buf);
ldpanel = lddp;
/* wait for the offdiagonal column */
magma_queue_wait_event( stream[d][stream1], event[d][1] );
} else {
dlpanel = dlA(d, 0, nb*j_local);
ldpanel = ldda;
}
/* update the panel */
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream1]);
trace_gpu_start( d, stream1, "gemm", "gemm" );
magma_cgemm(MagmaConjTrans, 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, stream1 );
}
d = (d+1)%num_gpus;
}
}
/* factor the diagonal */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] ); // wait for the diagonal
trace_cpu_start( 0, "getrf", "getrf" );
lapackf77_cpotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
trace_cpu_end( 0 );
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++ ) {
magma_setdevice(d);
if ( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d, 0, 0, buf);
ldpanel = lddp;
}
trace_gpu_start( d, stream0, "comm", "D to GPUs" );
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, ldpanel,
stream[d][stream0] );
trace_gpu_end( d, stream0 );
magma_event_record( event[d][2], stream[d][stream0] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
trace_gpu_start( id, stream0, "comm", "D to GPUs" );
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
stream[id][stream0] );
trace_gpu_end( id, stream0 );
}
/* 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]-nb*j_local2;
nb0 = min(nb, nb2 );
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream1]);
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for the diagonal
if ( j+jb < m && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead the column */
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if defined(PRECISION_d) && defined(CTRSM_WORK)
magmablas_ctrsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d][0], d_x[d][0] );
/*nb2 = n_local[d] - j_local2*nb;
magmablas_ctrsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d], d_x[d] ); */
#else
/*nb2 = n_local[d] - j_local2*nb;
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldda,
dlA(d, j, nb*j_local2), ldda);
*/
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream1 );
magma_event_record( event[d][3], stream[d][stream1] );
/* send the column to cpu */
if ( j+jb < m ) {
trace_gpu_start( d, stream0, "comm", "rows to CPU" );
magma_queue_wait_event( stream[d][stream0], event[d][3] ); // wait for lookahead
magma_cgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
stream[d][stream0] );
trace_gpu_end( d, stream0 );
magma_event_record( event[d][1], stream[d][stream0] );
}
/* update the remaining blocks */
nb2 = nb2 - nb0;
#if defined(PRECISION_d) && defined(CTRSM_WORK)
magmablas_ctrsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, 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_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda);
#endif
} else if ( nb2 > 0 ) {
/* update the entire trailing matrix */
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if defined(PRECISION_d) && defined(CTRSM_WORK)
magmablas_ctrsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream1 );
}
d = (d+1)%num_gpus;
}
} /* end of ctrsm */
} /* end of for j=1, .., n */
} else {
/* -------------------------------------------- */
/* Lower-triangular case */
/* Compute the Cholesky factorization A = L*L'. */
/* -------------------------------------------- */
#if defined(PRECISION_d) && defined(CTRSM_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++ ) {
magma_cmalloc( &d_dinvA[d][j], nb*nb );
magma_cmalloc( &d_x[d][j], nb*m );
cudaMemset(d_dinvA[d][j], 0, nb*nb*sizeof(magmaFloatComplex));
cudaMemset(d_x[d][j], 0, nb* m*sizeof(magmaFloatComplex));
}
}
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));
if ( j > 0 ) {
/* needed on pluto... */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] ); // 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 ) {
magma_setdevice(d);
/* wait for it on CPU */
magma_queue_wait_event( stream[d][stream0], event[id][1] );
/* send it to GPU */
magma_csetmatrix_async( jb, j,
Alo(j,0), lda,
dlPT(d,0,jb,buf), nb,
stream[d][stream0] );
magma_event_record( event[d][1], stream[d][stream0] );
}
d = (d+1)%num_gpus;
}
}
/* Update the current diagonal block */
magma_setdevice(id);
if ( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
magma_cherk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda);
magma_event_record( event[id][0], stream[id][stream1] );
}
/* send the diagonal to cpu */
magma_queue_wait_event( stream[id][stream0], event[id][0] ); // wait for syrk
magma_cgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
stream[id][stream0] );
/* 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, 0, jb, buf);
ldpanel = nb;
/* wait for offdiagonal row */
magma_queue_wait_event( stream[d][stream1], event[d][1] );
} else {
dlpanel = dlA(d, nb*j_local, 0);
ldpanel = ldda;
}
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream1]);
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, ldpanel,
c_one, dlA(d, nb0, j), ldda);
}
d = (d+1)%num_gpus;
}
}
/* factor the diagonal */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] );
lapackf77_cpotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
if (*info != 0) {
*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++ ) {
magma_setdevice(d);
if ( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, ldpanel,
stream[d][stream0] );
magma_event_record( event[d][2], stream[d][stream0] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
stream[id][stream0] );
}
/* 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 = 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);
magmablasSetKernelStream(stream[d][stream1]);
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for the diagonal
if ( j+jb < n && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead the column */
#if defined(PRECISION_d) && defined(CTRSM_WORK)
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
d_dinvA[d][0], d_x[d][0]);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
magma_event_record( event[d][3], stream[d][stream1] );
/* send the column to cpu */
if ( j+jb < n ) {
magma_queue_wait_event( stream[d][stream0], event[d][3] ); // wait for lookahead
magma_cgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
stream[d][stream0] );
magma_event_record( event[d][1], stream[d][stream0] );
}
/* update the remaining blocks */
nb2 = nb2 - nb0;
#if defined(PRECISION_d) && defined(CTRSM_WORK)
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, 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_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda);
#endif
} else if ( nb2 > 0 ) {
/* update the entire trailing matrix */
#if defined(PRECISION_d) && defined(CTRSM_WORK)
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
}
d = (d+1)%num_gpus;
}
}
}
} /* end of else not upper */
/* == finalize the trace == */
trace_finalize( "cpotrf.svg", "trace.css" );
/* clean up */
for( d=0; d < num_gpus; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
//magma_event_destroy( event0[d] );
//magma_event_destroy( event1[d] );
//magma_event_destroy( event2[d] );
//magma_event_destroy( event3[d] );
}
magma_setdevice(0);
return *info;
} /* magma_cpotrf_mgpu */
/**
Purpose
-------
CPOTRF computes the Cholesky factorization of a complex Hermitian
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 COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If uplo = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If uplo = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\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_cposv_comp
********************************************************************/
extern "C" magma_int_t
magma_cpotrf(
magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *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 magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const float d_one = 1.0;
const float 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;
magmaFloatComplex_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_cpotrf_nb( n );
if (nb <= 1 || nb >= n) {
lapackf77_cpotrf( 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_cpotrf_m( ngpu, uplo, n, A, lda, info );
}
if (MAGMA_SUCCESS != magma_cmalloc( &dA, n*ldda )) {
/* alloc failed so call the non-GPU-resident version */
return magma_cpotrf_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_csetmatrix_async( jb, n-j,
A(j, j), lda,
dA(j, j), ldda, queues[1] );
magma_cherk( 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_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, queues[0] );
if (j+jb < n) {
magma_cgemm( 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_cgetmatrix_async( j, jb,
dA(0, j), ldda,
A(0, j), lda, queues[0] );
lapackf77_cpotrf( MagmaUpperStr, &jb, A(j, j), &lda, info );
if (*info != 0) {
*info = *info + j;
break;
}
magma_csetmatrix_async( jb, jb,
A(j, j), lda,
dA(j, j), ldda, queues[0] );
magma_queue_sync( queues[0] );
if (j+jb < n) {
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, n-j-jb,
c_one, dA(j, j ), ldda,
dA(j, j+jb), ldda, queues[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.
jb = min( nb, n-j );
magma_csetmatrix_async( n-j, jb,
A(j, j), lda,
dA(j, j), ldda, queues[1] );
magma_cherk( 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_cgetmatrix_async( jb, jb,
dA(j,j), ldda,
A(j,j), lda, queues[0] );
if (j+jb < n) {
magma_cgemm( 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] );
}
magma_queue_sync( queues[0] );
// this could be on any queue; it isn't needed until exit.
magma_cgetmatrix_async( jb, j,
dA(j, 0), ldda,
A(j, 0), lda, queues[0] );
lapackf77_cpotrf( MagmaLowerStr, &jb, A(j, j), &lda, info );
if (*info != 0) {
*info = *info + j;
break;
}
magma_csetmatrix_async( jb, jb,
A(j, j), lda,
dA(j, j), ldda, queues[0] );
magma_queue_sync( queues[0] );
if (j+jb < n) {
magma_ctrsm( 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_cpotrf */
/**
Purpose
-------
CLAUUM 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 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 COPLEX_16 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';
if UPLO = MagmaLower, the lower triangle of A is overwritten with
the lower triangle of the product L' * 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_cposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_clauum(magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *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 i, ib;
magmaFloatComplex c_one = MAGMA_C_ONE;
float d_one = MAGMA_D_ONE;
magmaFloatComplex *dA;
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;
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
nb = magma_get_cpotrf_nb(n);
if (nb <= 1 || nb >= n)
lapackf77_clauum(uplo_, &n, A, &lda, info);
else {
if (upper) {
/* Compute the product U * U'. */
for (i=0; i < n; i += nb) {
ib=min(nb,n-i);
magma_csetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_csetmatrix_async( ib, (n-i-ib),
A(i,i+ib), lda,
dA(i,i+ib), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ctrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
lapackf77_clauum(MagmaUpperStr, &ib, A(i,i), &lda, info);
magma_csetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_cgemm( 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);
magma_queue_sync( stream[0] );
magma_cherk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
magma_cgetmatrix( i+ib, ib,
dA(0, i), ldda,
A(0, i), lda );
}
}
else {
/* Compute the product L' * L. */
for (i=0; i < n; i += nb) {
ib=min(nb,n-i);
magma_csetmatrix_async( ib, ib,
A(i,i), lda,
dA(i, i), ldda, stream[1] );
magma_csetmatrix_async( (n-i-ib), ib,
A(i+ib, i), lda,
dA(i+ib, i), ldda, stream[0] );
magma_queue_sync( stream[1] );
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
lapackf77_clauum(MagmaLowerStr, &ib, A(i,i), &lda, info);
magma_csetmatrix_async( ib, ib,
A(i, i), lda,
dA(i, i), ldda, stream[0] );
if (i+ib < n) {
magma_cgemm(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);
magma_queue_sync( stream[0] );
magma_cherk(MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
magma_cgetmatrix( ib, i+ib,
dA(i, 0), ldda,
A(i, 0), lda );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dA );
return *info;
}
/**
Purpose
-------
CPOTRF computes the Cholesky factorization of a complex Hermitian
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]
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 COMPLEX array of pointers on the GPU, dimension (ngpu)
On entry, the Hermitian 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_cposv_comp
********************************************************************/
extern "C" magma_int_t
magma_cpotrf_mgpu_right(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_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)
magmaFloatComplex z_one = MAGMA_C_MAKE( 1.0, 0.0 );
magmaFloatComplex mz_one = MAGMA_C_MAKE( -1.0, 0.0 );
float one = 1.0;
float m_one = -1.0;
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t j, nb, d, id, j_local, blkid, crosspoint, prevj, prevtrsmrows=0, nqueue = 5;
magmaFloatComplex *panel, *tmppanel0, *tmppanel1, *tmppanel, *tmpprevpanel;
magmaFloatComplex *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_cpotrf_nb(n);
ldpanel = ldda;
magma_setdevice(0);
if (MAGMA_SUCCESS != magma_cmalloc_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_cgetmatrix( n, n, dlA(0, 0, 0), ldda, panel, ldpanel);
lapackf77_cpotrf( uplo_, &n, panel, &ldpanel, info);
magma_csetmatrix( 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_cmalloc( &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_cgetmatrix_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_cgemm( MagmaNoTransStr, MagmaConjTransStr,
&rows, &nb, &nb,
&mz_one, tmpprevpanel(j), &ldpanel,
tmpprevpanel(j), &ldpanel,
&z_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_cpotrf(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_ctrsm(MagmaRightStr, MagmaLowerStr, MagmaConjTransStr, MagmaNonUnitStr,
&trsmrows, &nb,
&z_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_csetmatrix_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 CHERK_ON_DIAG
#ifdef CHERK_ON_DIAG
magma_cherk( MagmaLower, MagmaNoTrans,
nb, nb,
m_one, dlpanel, ldda,
one, dlA(d, j + nb, j_local2), ldda);
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows-nb, nb, nb,
mz_one, dlpanel+nb, ldda,
dlpanel, ldda,
z_one, dlA(d, j + nb +nb, j_local2), ldda);
#else
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows, nb, nb,
mz_one, dlpanel, ldda,
dlpanel, ldda,
z_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_cgetmatrix_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_cherk(MagmaLower, MagmaNoTrans, n - offset, nb,
// m_one, dlpanel, ldda,
// one, &d_lA[d][offset + offset*ldda], ldda );
#ifdef CHERK_ON_DIAG
magma_cherk_mgpu
#else
magma_cherk_mgpu2
#endif
(ngpu, MagmaLower, MagmaNoTrans,
nb, n - offset, nb,
m_one, dlpanels, ldda, 0,
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;
prevj = j;
#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_cgetmatrix(rows, rows, dlA(id, j, j_local), ldda, panel(j), ldpanel);
lapackf77_cpotrf(MagmaLowerStr, &rows, panel(j), &ldpanel, info);
magma_csetmatrix(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_cpotrf_mgpu_right */
