/**
Purpose
-------
CTRTRI computes the inverse of a real upper or lower triangular
matrix dA.
This is the Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: A is upper triangular;
- = MagmaLower: A is lower triangular.
@param[in]
diag magma_diag_t
- = MagmaNonUnit: A is non-unit triangular;
- = MagmaUnit: A is unit triangular.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array ON THE GPU, dimension (LDDA,N)
On entry, the triangular matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of the array dA contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of the array dA contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = MagmaUnit, the
diagonal elements of A are also not referenced and are
assumed to be 1.
On exit, the (triangular) inverse of the original matrix, in
the same storage format.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, dA(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
(Singularity check is currently disabled.)
@ingroup magma_cgesv_aux
********************************************************************/
extern "C" magma_int_t
magma_ctrtri_gpu(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaFloatComplex_ptr 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 );
const char* diag_ = lapack_diag_const( diag );
magma_int_t nb, nn, j, jb;
//magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *work;
int upper = (uplo == MagmaUpper);
int nounit = (diag == MagmaNonUnit);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (! nounit && diag != MagmaUnit)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Check for singularity if non-unit */
/* cannot do here with matrix dA on GPU -- need kernel */
/*
if (nounit) {
for (j=0; j < n; ++j) {
if ( MAGMA_C_EQUAL( *dA(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
*/
/* Determine the block size for this environment */
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_ctrtri( uplo_, diag_, &n, work, &n, info );
magma_csetmatrix( n, n, work, n, dA, ldda );
}
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
/* Compute rows 1:j-1 of current block column */
magma_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j), ldda );
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j), ldda );
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaUpperStr, diag_, &jb, work, &jb, info );
magma_csetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[0] );
}
}
else {
/* Compute inverse of lower triangular matrix */
nn = ((n-1)/nb)*nb+1;
for (j=nn-1; j >= 0; j -= nb) {
jb = min(nb,(n-j));
if ((j+jb) < n) {
/* Compute rows j+jb:n of current block column */
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
}
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaLowerStr, diag_, &jb, work, &jb, info );
magma_csetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[0] );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
}
/**
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" magma_int_t
magma_cpotrf_msub(
magma_int_t num_subs, magma_int_t num_gpus, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr *d_lA, size_t dA_offset,
magma_int_t ldda,
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 err;
magma_int_t j, nb, d, lddp, h;
magmaFloatComplex *work;
magmaFloatComplex_ptr dwork[MagmaMaxGPUs];
*info = 0;
nb = magma_get_cpotrf_nb(n);
if ( uplo != MagmaUpper && uplo != MagmaLower ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (uplo != MagmaUpper) {
lddp = nb*(n/(nb*tot_subs));
if( n%(nb*tot_subs) != 0 ) lddp+=min(nb,n-tot_subs*lddp);
if( ldda < lddp ) *info = -4;
} else if( ldda < n ) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (num_gpus == 1 && ((nb <= 1) || (nb >= n)) ) {
/* Use unblocked code. */
err = magma_cmalloc_cpu( &work, n*nb );
if (err != MAGMA_SUCCESS) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( n, n, d_lA[0], 0, ldda, work, n, queues[0] );
lapackf77_cpotrf(lapack_uplo_const(uplo), &n, work, &n, info);
magma_csetmatrix( n, n, work, n, d_lA[0], 0, ldda, queues[0] );
magma_free_cpu( work );
} else {
lddp = 32*((n+31)/32);
for (d=0; d<num_gpus; d++) {
if (MAGMA_SUCCESS != magma_cmalloc( &dwork[d], num_gpus*nb*lddp )) {
for( j=0; j<d; j++ ) magma_free( dwork[j] );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
h = 1; //num_gpus; //(n+nb-1)/nb;
#ifdef USE_PINNED_CLMEMORY
cl_mem buffer = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(magmaFloatComplex)*n*nb*h, NULL, NULL);
for (d=0; d<num_gpus; d++) {
work = (magmaFloatComplex*)clEnqueueMapBuffer(queues[2*d], buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
sizeof(magmaFloatComplex)*n*nb*h, 0, NULL, NULL, NULL);
}
#else
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &work, n*nb*h )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
if (uplo == MagmaUpper) {
/* with two queues for each device */
magma_cpotrf2_msub(num_subs, num_gpus, uplo, n, n, 0, 0, nb, d_lA, 0, ldda,
dwork, lddp, work, n, h, queues, info);
//magma_cpotrf3_msub(num_subs, num_gpus, uplo, n, n, 0, 0, nb, d_lA, 0, ldda,
// dwork, lddp, work, n, h, queues, info);
/* with three streams */
//magma_cpotrf3_msub(num_gpus, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, n,
// h, stream, event, info);
} else {
/* with two queues for each device */
magma_cpotrf2_msub(num_subs, num_gpus, uplo, n, n, 0, 0, nb, d_lA, 0, ldda,
dwork, lddp, work, nb*h, h, queues, info);
//magma_cpotrf3_msub(num_subs, num_gpus, uplo, n, n, 0, 0, nb, d_lA, 0, ldda,
// dwork, lddp, work, nb*h, h, queues, info);
//magma_cpotrf4_msub(num_subs, num_gpus, uplo, n, n, 0, 0, nb, d_lA, 0, ldda,
// dwork, lddp, work, nb*h, h, queues, info);
/* with three streams */
//magma_cpotrf3_msub(num_gpus, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, nb*h,
// h, stream, event, info);
}
/* clean up */
for (d=0; d<num_gpus; d++) magma_free( dwork[d] );
#ifdef USE_PINNED_CLMEMORY
for (d=0; d<num_gpus; d++) {
clEnqueueUnmapMemObject(queues[2*d], buffer, work, 0, NULL, NULL);
}
clReleaseMemObject( buffer );
#else
magma_free_cpu( work );
#endif
} /* end of not lapack */
return *info;
} /* magma_cpotrf_msub */
extern "C" magma_int_t
magma_cpotrf_mgpu(magma_int_t num_gpus, char uplo, magma_int_t n,
magmaFloatComplex **d_lA, magma_int_t ldda, magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 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.
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, nb, d, lddp, h;
char uplo_[2] = {uplo, 0};
magmaFloatComplex *work;
int upper = lapackf77_lsame(uplo_, "U");
magmaFloatComplex *dwork[MagmaMaxGPUs];
magma_queue_t stream[MagmaMaxGPUs][3];
magma_event_t event[MagmaMaxGPUs][5];
*info = 0;
nb = magma_get_cpotrf_nb(n);
if ( (! upper) && (! lapackf77_lsame(uplo_, "L")) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (!upper) {
lddp = nb*(n/(nb*num_gpus));
if( n%(nb*num_gpus) != 0 ) lddp+=min(nb,n-num_gpus*lddp);
if( ldda < lddp ) *info = -4;
} else if( ldda < n ) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (num_gpus == 1 && ((nb <= 1) || (nb >= n)) ) {
/* Use unblocked code. */
magma_setdevice(0);
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, n*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( n, n, d_lA[0], ldda, work, n );
lapackf77_cpotrf(uplo_, &n, work, &n, info);
magma_csetmatrix( n, n, work, n, d_lA[0], ldda );
magma_free_pinned( work );
}
else {
lddp = nb*((n+nb-1)/nb);
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_cmalloc( &dwork[d], num_gpus*nb*lddp )) {
for( j=0; j<d; j++ ) {
magma_setdevice(j);
magma_free( dwork[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
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_setdevice(0);
h = 1; //num_gpus; //(n+nb-1)/nb;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, n*nb*h )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (upper) {
/* with two streams */
//magma_cpotrf2_mgpu(num_gpus, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, n,
// h, stream, event, info);
/* with three streams */
magma_cpotrf3_mgpu(num_gpus, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, n,
h, stream, event, info);
} else {
/* with two streams */
//magma_cpotrf2_mgpu(num_gpus, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, nb*h,
// h, stream, event, info);
/* with three streams */
magma_cpotrf3_mgpu(num_gpus, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, nb*h,
h, stream, event, info);
}
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j<3; j++ ) {
magma_queue_sync( stream[d][j] );
magma_queue_destroy( stream[d][j] );
}
magmablasSetKernelStream(NULL);
for( j=0; j<5; j++ )
magma_event_destroy( event[d][j] );
magma_free( dwork[d] );
}
magma_setdevice(0);
magma_free_pinned( work );
} /* end of not lapack */
return *info;
} /* magma_cpotrf_mgpu */
extern "C" magma_int_t
magma_ctrtri_gpu(
magma_uplo_t uplo, magma_diag_t diag, magma_int_t n,
magmaFloatComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_queue_t queues[2],
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
=======
CTRTRI computes the inverse of a real upper or lower triangular
matrix dA.
This is the Level 3 BLAS version of the algorithm.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': A is upper triangular;
= 'L': A is lower triangular.
DIAG (input) CHARACTER*1
= 'N': A is non-unit triangular;
= 'U': A is unit triangular.
N (input) INTEGER
The order of the matrix A. N >= 0.
dA (input/output) COMPLEX array ON THE GPU, dimension (LDDA,N)
On entry, the triangular matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of the array dA contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of the array dA contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = 'U', the
diagonal elements of A are also not referenced and are
assumed to be 1.
On exit, the (triangular) inverse of the original matrix, in
the same storage format.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, dA(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
(Singularity check is currently disabled.)
===================================================================== */
/* Local variables */
magma_int_t nb, nn, j, jb;
//magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *work;
int upper = (uplo == MagmaUpper);
int nounit = (diag == MagmaNonUnit);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (! nounit && diag != MagmaUnit)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Check for singularity if non-unit */
/* cannot do here with matrix dA on GPU -- need kernel */
/*
if (nounit) {
for (j=0; j < n; ++j) {
if ( MAGMA_C_EQUAL( *dA(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
*/
/* Determine the block size for this environment */
nb = magma_get_cpotrf_nb(n);
/* Create Queues */
//magma_queue_t queues[2];
//magma_device_t device[MagmaMaxGPUs];
//magma_int_t num = 0;
//magma_int_t err;
//
//err = magma_getdevices( device, MagmaMaxGPUs, &num );
//if ( err != 0 || num < 1 ) {
// fprintf( stderr, "magma_getdevices failed: %d\n", err );
// exit(-1);
//}
//err = magma_queue_create( device[0], &queues[0] );
//if ( err != 0 ) {
// fprintf( stderr, "magma_queue_create 0 failed: %d\n", err );
// exit(-1);
//}
//err = magma_queue_create( device[0], &queues[1] );
//if ( err != 0 ) {
// fprintf( stderr, "magma_queue_create 1 failed: %d\n", err );
// exit(-1);
//}
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (nb <= 1 || nb >= n) {
magma_cgetmatrix( n, n, dA, dA_offset, ldda, work, n, queues[0] );
lapackf77_ctrtri( lapack_const(uplo), lapack_const(diag), &n, work, &n, info );
magma_csetmatrix( n, n, work, n, dA, dA_offset, ldda, queues[0] );
}
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
/* Compute rows 1:j-1 of current block column */
magma_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j), ldda,
queues[0] );
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j), ldda,
queues[0] );
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaUpperStr, lapack_const(diag), &jb, work, &jb, info );
/*
magma_csetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_csetmatrix( jb, jb,
work, jb,
dA(j, j), ldda, queues[0] );
}
}
else {
/* Compute inverse of lower triangular matrix */
nn = ((n-1)/nb)*nb+1;
for(j=nn-1; j >= 0; j -= nb) {
jb = min(nb,(n-j));
if((j+jb) < n) {
/* Compute rows j+jb:n of current block column */
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda,
queues[0] );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda,
queues[0] );
}
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[1], NULL );
magma_queue_sync( queues[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri( MagmaLowerStr, lapack_const(diag), &jb, work, &jb, info );
/*
magma_csetmatrix_async( jb, jb,
work, 0, jb,
dA(j, j), ldda, queues[0], NULL );
*/
magma_csetmatrix( jb, jb,
work, jb,
dA(j, j), ldda, queues[0] );
}
}
}
//magma_queue_destroy( queues[0] );
//magma_queue_destroy( queues[1] );
magma_free_cpu( work );
return *info;
}
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 */
/**
Purpose
-------
CTRTRI computes the inverse of a real upper or lower triangular
matrix A.
This is the Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: A is upper triangular;
- = MagmaLower: A is lower triangular.
@param[in]
diag magma_diag_t
- = MagmaNonUnit: A is non-unit triangular;
- = MagmaUnit: A is unit triangular.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the triangular matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of the array A contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of the array A contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = MagmaUnit, the
diagonal elements of A are also not referenced and are
assumed to be 1.
On exit, the (triangular) inverse of the original matrix, in
the same storage format.
@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
- > 0: if INFO = i, A(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_ctrtri(
magma_uplo_t uplo, magma_diag_t diag, 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_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const char* uplo_ = lapack_uplo_const( uplo );
const char* diag_ = lapack_diag_const( diag );
// Local variables
magma_int_t ldda, nb, nn, j, jb;
magmaFloatComplex_ptr dA;
bool upper = (uplo == MagmaUpper);
bool nounit = (diag == MagmaNonUnit);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (! nounit && diag != MagmaUnit)
*info = -2;
else if (n < 0)
*info = -3;
else if (lda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// Quick return
if ( n == 0 )
return *info;
// Check for singularity if non-unit
if (nounit) {
for (j=0; j < n; ++j) {
if ( MAGMA_C_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
// Determine the block size for this environment
nb = magma_get_cpotrf_nb( n );
ldda = magma_roundup( n, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
//magma_queue_create( cdev, &queues[1] ); // unused
if (nb <= 1 || nb >= n) {
lapackf77_ctrtri( uplo_, diag_, &n, A, &lda, info );
}
else if (upper) {
// Compute inverse of upper triangular matrix
for (j=0; j < n; j += nb) {
jb = min( nb, n-j );
if (j > 0) {
// Send current block column (with diagonal) to device
// This must finish before trtri below
magma_csetmatrix( j+jb, jb,
A(0,j), lda,
dA(0,j), ldda, queues[0] );
// Compute rows 0:j of current block column
magma_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, diag, j, jb, c_one,
dA(0,0), ldda,
dA(0,j), ldda, queues[0] );
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, diag, j, jb, c_neg_one,
dA(j,j), ldda,
dA(0,j), ldda, queues[0] );
// Get above diagonal from device
// TODO: could be on another queue, after trmm/trsm finish
magma_cgetmatrix_async( j, jb,
dA(0,j), ldda,
A(0,j), lda, queues[0] );
}
// Compute inverse of current diagonal block
// TODO: problem if diagonal has not finished sending yet?
lapackf77_ctrtri( MagmaUpperStr, diag_, &jb, A(j,j), &lda, info );
if (j+jb < n) {
// Send inverted diagonal block to device
magma_csetmatrix( jb, jb,
A(j,j), lda,
dA(j,j), ldda, queues[0] );
}
}
}
else {
// Compute inverse of lower triangular matrix
nn = ((n-1)/nb)*nb;
for (j=nn; j >= 0; j -= nb) {
jb = min( nb, n-j );
if (j+jb < n) {
// Send current block row (with diagonal) to device
// This must finish before trtri below
magma_csetmatrix( n-j, jb,
A(j,j), lda,
dA(j,j), ldda, queues[0] );
// Compute rows j+jb:n of current block column
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, diag, n-j-jb, jb, c_one,
dA(j+jb,j+jb), ldda,
dA(j+jb,j), ldda, queues[0] );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, diag, n-j-jb, jb, c_neg_one,
dA(j,j), ldda,
dA(j+jb,j), ldda, queues[0] );
// Get below diagonal block from device
magma_cgetmatrix_async( n-j-jb, jb,
dA(j+jb,j), ldda,
A(j+jb,j), lda, queues[0] );
}
// Compute inverse of current diagonal block
lapackf77_ctrtri( MagmaLowerStr, diag_, &jb, A(j,j), &lda, info );
if (j > 0) {
// Send inverted diagonal block to device
magma_csetmatrix( jb, jb,
A(j,j), lda,
dA(j,j), ldda, queues[0] );
}
}
}
magma_queue_destroy( queues[0] );
//magma_queue_destroy( queues[1] ); // unused
magma_free( dA );
return *info;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
magmaFloatComplex *h_R = NULL, *h_P = NULL;
magmaFloatComplex_ptr d_lA[MagmaMaxSubs * MagmaMaxGPUs];
magma_int_t N = 0, n2, lda, ldda;
magma_int_t size[10] =
{ 1000, 2000, 3000, 4000, 5000, 6000, 7000, 8000, 9000, 10000 };
magma_int_t i, j, k, check = 0, info;
magmaFloatComplex mz_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t num_gpus0 = 1, num_gpus, num_subs0 = 1, num_subs, tot_subs, flag = 0;
int nb, n_local, nk;
magma_uplo_t uplo = MagmaLower;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i]) == 0){
N = atoi(argv[++i]);
if (N > 0) {
size[0] = size[9] = N;
flag = 1;
}
}
if(strcmp("-NGPU", argv[i]) == 0)
num_gpus0 = atoi(argv[++i]);
if(strcmp("-NSUB", argv[i]) == 0)
num_subs0 = atoi(argv[++i]);
if(strcmp("-UPLO", argv[i]) == 0)
uplo = (strcmp("L", argv[++i]) == 0 ? MagmaLower : MagmaUpper);
if(strcmp("-check", argv[i]) == 0)
check = 1;
}
}
/* Initialize */
magma_queue_t queues[2*MagmaMaxGPUs];
magma_device_t devices[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( devices, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
for(i=0;i<num_gpus0;i++){
err = magma_queue_create( devices[i], &queues[2*i] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( devices[i], &queues[2*i+1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
}
printf("\nUsing %d GPUs:\n", num_gpus0);
printf(" testing_cpotrf_msub -N %d -NGPU %d -NSUB %d -UPLO %c %s\n\n", size[0], num_gpus0,num_subs0,
(uplo == MagmaLower ? 'L' : 'U'),(check == 1 ? "-check" : " "));
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R_magma-R_lapack||_F / ||R_lapack||_F\n");
printf("========================================================================================\n");
for(i=0; i<10; i++){
N = size[i];
lda = N;
n2 = lda*N;
gflops = FLOPS_CPOTRF( N ) / 1e9;;
nb = magma_get_cpotrf_nb(N);
if (num_subs0*num_gpus0 > N/nb) {
num_gpus = N/nb;
num_subs = 1;
if(N%nb != 0) num_gpus ++;
printf("too many GPUs for the matrix size, using %d GPUs\n", (int)num_gpus);
} else {
num_gpus = num_gpus0;
num_subs = num_subs0;
}
tot_subs = num_subs * num_gpus;
/* Allocate host memory for the matrix */
#ifdef USE_PINNED_CLMEMORY
cl_mem buffer1 = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, n2*sizeof(magmaFloatComplex), NULL, NULL);
cl_mem buffer2 = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, lda*nb*sizeof(magmaFloatComplex), NULL, NULL);
for (k=0; k<num_gpus; k++) {
h_R = (magmaFloatComplex*)clEnqueueMapBuffer(queues[2*k], buffer1, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
n2*sizeof(magmaFloatComplex), 0, NULL, NULL, NULL);
h_P = (magmaFloatComplex*)clEnqueueMapBuffer(queues[2*k], buffer2, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
lda*nb*sizeof(magmaFloatComplex), 0, NULL, NULL, NULL);
}
#else
TESTING_MALLOC_PIN( h_P, magmaFloatComplex, lda*nb );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
#endif
/* Initialize the matrix */
init_matrix( N, h_R, lda );
/* Allocate GPU memory */
if (uplo == MagmaUpper) {
ldda = ((N+nb-1)/nb)*nb;
n_local = ((N+nb*tot_subs-1)/(nb*tot_subs))*nb;
} else {
ldda = ((N+nb*tot_subs-1)/(nb*tot_subs))*nb;
n_local = ((N+nb-1)/nb)*nb;
}
for (j=0; j<tot_subs; j++) {
TESTING_MALLOC_DEV( d_lA[j], magmaFloatComplex, n_local*ldda );
}
/* Warm up to measure the performance */
/* distribute matrix to gpus */
if (uplo == MagmaUpper) {
for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_csetmatrix(j+nk, nk,
&h_R[j*lda], 0, lda,
d_lA[k], j/(nb*tot_subs)*nb*ldda, ldda,
queues[2*(k%num_gpus)]);
}
} else {
for (j=0; j<N; j+=nb) {
nk = min(nb, N-j);
for (int kk = 0; kk<tot_subs; kk++) {
int mk = 0;
for (int ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
int mii = min(nb, N-ii);
lapackf77_clacpy( MagmaFullStr, &mii, &nk, &h_R[ii+j*lda], &lda, &h_P[mk], &lda );
mk += mii;
}
k = ((j+kk*nb)/nb)%tot_subs;
if (mk > 0 && nk > 0) {
magma_csetmatrix(mk, nk,
h_P, 0, lda,
d_lA[k], j*ldda+(j+kk*nb)/(nb*tot_subs)*nb, ldda,
queues[2*(k%num_gpus)]);
}
}
}
/*for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_csetmatrix(nk, j+nk, &h_R[j], 0, lda,
d_lA[k], j/(nb*tot_subs)*nb, ldda,
queues[2*(k%num_gpus)]);
}*/
}
magma_cpotrf_msub( num_subs, num_gpus, uplo, N, d_lA, 0, ldda, &info, queues );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
/* distribute matrix to gpus */
if (uplo == MagmaUpper) {
for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_csetmatrix(j+nk, nk,
&h_R[j*lda], 0, lda,
d_lA[k], j/(nb*tot_subs)*nb*ldda, ldda,
queues[2*(k%num_gpus)]);
}
} else {
for (j=0; j<N; j+=nb) {
nk = min(nb, N-j);
for (int kk = 0; kk<tot_subs; kk++) {
int mk = 0;
for (int ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
int mii = min(nb, N-ii);
lapackf77_clacpy( MagmaFullStr, &mii, &nk, &h_R[ii+j*lda], &lda, &h_P[mk], &lda );
mk += mii;
}
k = ((j+kk*nb)/nb)%tot_subs;
if (mk > 0 && nk > 0) {
magma_csetmatrix(mk, nk,
h_P, 0, lda,
d_lA[k], j*ldda+(j+kk*nb)/(nb*tot_subs)*nb, ldda,
queues[2*(k%num_gpus)]);
}
}
}
/*for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_csetmatrix(nk, j+nk, &h_R[j], 0, lda,
d_lA[k], (j/(nb*tot_subs)*nb), ldda,
queues[2*(k%num_gpus)]);
}*/
}
gpu_time = magma_wtime();
magma_cpotrf_msub( num_subs, num_gpus, uplo, N, d_lA, 0, ldda, &info, queues );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf( "magma_cpotrf had error %d.\n", info );
/* gather matrix from gpus */
if (uplo==MagmaUpper) {
for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_cgetmatrix(j+nk, nk,
d_lA[k], j/(nb*tot_subs)*nb*ldda, ldda,
&h_R[j*lda], 0, lda, queues[2*(k%num_gpus)]);
}
} else {
for (j=0; j<N; j+=nb) {
nk = min(nb, N-j);
for (int kk = 0; kk<tot_subs; kk++) {
k = ((j+kk*nb)/nb)%tot_subs;
int mk = 0;
mk = 0;
for (int ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
mk += min(nb, N-ii);
}
if (mk > 0 && nk > 0) {
magma_cgetmatrix(mk, nk,
d_lA[k], j*ldda+(j+kk*nb)/(nb*tot_subs)*nb, ldda,
h_P, 0, lda,
queues[2*(k%num_gpus)]);
}
mk = 0;
for (int ii=j+kk*nb; ii<N; ii+=nb*tot_subs) {
int mii = min(nb, N-ii);
lapackf77_clacpy( MagmaFullStr, &mii, &nk, &h_P[mk], &lda, &h_R[ii+j*lda], &lda );
mk += mii;
}
}
}
/*for (j=0; j<N; j+=nb) {
k = (j/nb)%tot_subs;
nk = min(nb, N-j);
magma_cgetmatrix( nk, j+nk,
d_lA[k], (j/(nb*tot_subs)*nb), ldda,
&h_R[j], 0, lda, queues[2*(k%num_gpus)] );
}*/
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if (check == 1) {
float work[1], matnorm, diffnorm;
magmaFloatComplex *h_A;
TESTING_MALLOC_PIN( h_A, magmaFloatComplex, n2 );
init_matrix( N, h_A, lda );
cpu_time = magma_wtime();
if (uplo == MagmaLower) {
lapackf77_cpotrf( MagmaLowerStr, &N, h_A, &lda, &info );
} else {
lapackf77_cpotrf( MagmaUpperStr, &N, h_A, &lda, &info );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf( "lapackf77_cpotrf had error %d.\n", info );
/* =====================================================================
Check the result compared to LAPACK
|R_magma - R_lapack| / |R_lapack|
=================================================================== */
matnorm = lapackf77_clange("f", &N, &N, h_A, &lda, work);
blasf77_caxpy(&n2, &mz_one, h_A, &ione, h_R, &ione);
diffnorm = lapackf77_clange("f", &N, &N, h_R, &lda, work);
printf( "%5d %6.2f (%6.2f) %6.2f (%6.2f) %e\n",
N, cpu_perf, cpu_time, gpu_perf, gpu_time, diffnorm / matnorm );
TESTING_FREE_PIN( h_A );
} else {
printf( "%5d - - (- -) %6.2f (%6.2f) - -\n",
N, gpu_perf, gpu_time );
}
// free memory
#ifdef USE_PINNED_CLMEMORY
for (k=0; k<num_gpus; k++) {
clEnqueueUnmapMemObject(queues[2*k], buffer1, h_R, 0, NULL, NULL);
clEnqueueUnmapMemObject(queues[2*k], buffer2, h_P, 0, NULL, NULL);
}
clReleaseMemObject(buffer1);
clReleaseMemObject(buffer2);
#else
TESTING_FREE_PIN( h_P );
TESTING_FREE_PIN( h_R );
#endif
for (j=0; j<tot_subs; j++) {
TESTING_FREE_DEV( d_lA[j] );
}
if (flag != 0)
break;
}
/* clean up */
for (i=0; i<num_gpus; i++) {
magma_queue_destroy( queues[2*i] );
magma_queue_destroy( queues[2*i+1] );
}
magma_finalize();
return 0;
}
extern "C" magma_int_t
magma_ctrtri(char uplo, char diag, magma_int_t n,
cuFloatComplex *a, magma_int_t lda, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
CTRTRI computes the inverse of a real upper or lower triangular
matrix A.
This is the Level 3 BLAS version of the algorithm.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': A is upper triangular;
= 'L': A is lower triangular.
DIAG (input) CHARACTER*1
= 'N': A is non-unit triangular;
= 'U': A is unit triangular.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the triangular matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of the array A contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of the array A contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = 'U', the
diagonal elements of A are also not referenced and are
assumed to be 1.
On exit, the (triangular) inverse of the original matrix, in
the same storage format.
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
> 0: if INFO = i, A(i,i) is exactly zero. The triangular
matrix is singular and its inverse cannot be computed.
===================================================================== */
/* Local variables */
char uplo_[2] = {uplo, 0};
char diag_[2] = {diag, 0};
magma_int_t ldda, nb, nn, j, jb;
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
cuFloatComplex *work;
int upper = lapackf77_lsame(uplo_, "U");
int nounit = lapackf77_lsame(diag_, "N");
*info = 0;
if ((! upper) && (! lapackf77_lsame(uplo_, "L")))
*info = -1;
else if ((! nounit) && (! lapackf77_lsame(diag_, "U")))
*info = -2;
else if (n < 0)
*info = -3;
else if (lda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
/* Check for singularity if non-unit */
if (nounit) {
for ( j=0; j<n; ++j ) {
if ( MAGMA_C_EQUAL( *A(j,j), c_zero )) {
*info = j+1; // Fortran index
return *info;
}
}
}
/* Determine the block size for this environment */
nb = magma_get_cpotrf_nb(n);
ldda = ((n+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &work, (n)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
cudaStream_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n)
lapackf77_ctrtri(uplo_, diag_, &n, a, &lda, info);
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (j=0; j<n; j=j+nb) {
jb = min(nb, (n-j));
magma_csetmatrix( jb, (n-j),
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows 1:j-1 of current block column */
magma_ctrmm( MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_ctrsm( MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
//cublasGetMatrix(j ,jb, sizeof( cuFloatComplex),
//dA(0, j), ldda, A(0, j), lda);
magma_cgetmatrix_async( jb, jb,
dA(j, j), ldda,
A(j, j), lda, stream[1] );
magma_cgetmatrix_async( j, jb,
dA(0, j), ldda,
A(0, j), lda, stream[0] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ctrtri(MagmaUpperStr, diag_, &jb, A(j,j), &lda, info);
magma_csetmatrix( jb, jb,
A(j, j), lda,
dA(j, j), ldda );
}
}
else {
/* Compute inverse of lower triangular matrix */
nn=((n-1)/nb)*nb+1;
for(j=nn-1; j>=0; j=j-nb) {
jb=min(nb,(n-j));
if((j+jb) < n) {
magma_csetmatrix( (n-j), jb,
A(j, j), lda,
dA(j, j), ldda );
/* Compute rows j+jb:n of current block column */
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda );
magma_ctrsm( MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda );
//cublasGetMatrix((n-j), jb, sizeof( cuFloatComplex),dA(j, j), ldda, A(j, j), lda);
magma_cgetmatrix_async( n-j-jb, jb,
dA(j+jb, j), ldda,
A(j+jb, j), lda, stream[1] );
magma_cgetmatrix_async( jb, jb,
dA(j,j), ldda,
A(j,j), lda, stream[0] );
magma_queue_sync( stream[0] );
}
/* Compute inverse of current diagonal block */
lapackf77_ctrtri(MagmaLowerStr, diag_, &jb, A(j,j), &lda, info);
magma_csetmatrix( jb, jb,
A(j, j), lda,
dA(j, j), ldda );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( work );
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]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
d_lA COMPLEX array of pointers on the GPU, dimension (ngpu)
On entry, the Hermitian matrix dA distributed over GPUs
(d_lA[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 d_lA. 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(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr d_lA[], magma_int_t ldda,
magma_int_t *info)
{
magma_int_t j, nb, d, lddp, h;
const char* uplo_ = lapack_uplo_const( uplo );
magmaFloatComplex *work;
bool upper = (uplo == MagmaUpper);
magmaFloatComplex *dwork[MagmaMaxGPUs];
magma_queue_t queues[MagmaMaxGPUs][3];
magma_event_t event[MagmaMaxGPUs][5];
*info = 0;
nb = magma_get_cpotrf_nb(n);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (!upper) {
lddp = nb*(n/(nb*ngpu));
if ( n%(nb*ngpu) != 0 ) lddp += min(nb, n-ngpu*lddp);
if ( ldda < lddp ) *info = -4;
} else if ( ldda < n ) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
if (ngpu == 1 && ((nb <= 1) || (nb >= n)) ) {
/* Use unblocked code. */
magma_setdevice(0);
magma_queue_create( 0, &queues[0][0] );
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, n*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( n, n, d_lA[0], ldda, work, n, queues[0][0] );
lapackf77_cpotrf(uplo_, &n, work, &n, info);
magma_csetmatrix( n, n, work, n, d_lA[0], ldda, queues[0][0] );
magma_free_pinned( work );
magma_queue_destroy( queues[0][0] );
}
else {
lddp = magma_roundup( n, nb );
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_cmalloc( &dwork[d], ngpu*nb*lddp )) {
for( j=0; j < d; j++ ) {
magma_setdevice(j);
magma_free( dwork[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
for( j=0; j < 3; j++ ) {
magma_queue_create( d, &queues[d][j] );
}
for( j=0; j < 5; j++ ) {
magma_event_create( &event[d][j] );
}
}
magma_setdevice(0);
h = 1; //ngpu; //magma_ceildiv( n, nb );
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, n*nb*h )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (upper) {
/* with three queues */
magma_cpotrf3_mgpu(ngpu, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, n,
h, queues, event, info);
} else {
/* with three queues */
magma_cpotrf3_mgpu(ngpu, uplo, n, n, 0, 0, nb, d_lA, ldda, dwork, lddp, work, nb*h,
h, queues, event, info);
}
/* clean up */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
for( j=0; j < 3; j++ ) {
magma_queue_sync( queues[d][j] );
magma_queue_destroy( queues[d][j] );
}
for( j=0; j < 5; j++ )
magma_event_destroy( event[d][j] );
magma_free( dwork[d] );
}
magma_free_pinned( work );
} /* end of not lapack */
magma_setdevice( orig_dev );
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 */
