extern "C" magma_int_t
magma_cgetrf2_piv(magma_int_t m, magma_int_t n, magma_int_t start, magma_int_t end,
magmaFloatComplex *A, magma_int_t lda, magma_int_t *ipiv, magma_int_t *info)
{
magma_int_t I, k1, k2, nb, incx, minmn;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0)
return *info;
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* initialize nb */
nb = magma_get_cgetrf_nb(m);
minmn = min( end, min(m,n) );
for( I=start; I < end-nb; I += nb ) {
incx = 1;
k1 = 1+I+nb;
k2 = minmn;
lapackf77_claswp(&nb, A(0,I), &lda, &k1, &k2, ipiv, &incx);
}
return *info;
} /* magma_cgetrf_piv */
/**
Purpose
-------
CGETRF_NOPIV_GPU computes an LU factorization of a general M-by-N
matrix A without any pivoting.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_nopiv_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define dA(i,j) (dA + (i)*nb + (j)*nb*ldda)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, mindim;
magma_int_t i, rows, s, lddwork;
magmaFloatComplex *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
magma_cgetrf_nopiv( m, n, work, m, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( i=0; i < s; i++ ) {
// download i-th panel
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dA(i,i), ldda, work, lddwork, stream[0] );
if ( i > 0 ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n - (i+1)*nb,
c_one, dA(i-1,i-1), ldda,
dA(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-i*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i, i-1), ldda, dA(i-1,i+1), ldda,
c_one, dA(i, i+1), ldda );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
magma_cgetrf_nopiv( rows, nb, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dA(i, i), ldda, stream[0] );
magma_queue_sync( stream[0] );
// do the small non-parallel computations
if ( s > (i+1) ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n-s*nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
magma_cgetmatrix( rows, nb0, dA(s,s), ldda, work, lddwork );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
magma_cgetrf_nopiv( rows, nb0, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dA(s,s), ldda );
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb0, n-s*nb-nb0,
c_one, dA(s,s), ldda,
dA(s,s)+nb0, ldda);
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_cgetrf_nopiv_gpu */
/**
Purpose
-------
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. 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 = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf(magma_int_t m, magma_int_t n, magmaFloatComplex *A, magma_int_t lda,
magma_int_t *ipiv, magma_int_t *info)
{
#define dAT(i,j) (dAT + (i)*nb*ldda + (j)*nb)
magmaFloatComplex *dAT, *dA, *da, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
nb = magma_get_cgetrf_nb(m);
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_cgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, maxdim;
magma_int_t i, rows, cols, s = min(m, n)/nb;
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
/* set number of GPUs */
magma_int_t num_gpus = magma_num_gpus();
if ( num_gpus > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
/* explicitly checking the memory requirement */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
int h = 1+(2+num_gpus), num_gpus2 = num_gpus;
int NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
if ( num_gpus > ceil((float)NB/nb) ) {
num_gpus2 = (int)ceil((float)NB/nb);
h = 1+(2+num_gpus2);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
}
if ( num_gpus2*NB < n ) {
/* require too much memory, so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
ldda = maxn;
work = A;
if (maxdim*maxdim < 2*maxm*maxn) {
// if close to square, allocate square matrix and transpose in-place
if (MAGMA_SUCCESS != magma_cmalloc( &dA, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
ldda = maxdim;
magma_csetmatrix( m, n, A, lda, da, ldda );
dAT = da;
magmablas_ctranspose_inplace( ldda, dAT, ldda );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
magma_csetmatrix( m, n, A, lda, da, maxm );
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dA );
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
magmablas_ctranspose( m, n, da, maxm, dAT, ldda );
}
lapackf77_cgetrf( &m, &nb, work, &lda, ipiv, &iinfo);
/* 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;
for( i = 0; i < s; i++ ) {
// download i-th panel
cols = maxm - i*nb;
if (i > 0) {
// download i-th panel
magmablas_ctranspose( nb, cols, dAT(i,i), ldda, dA, cols );
// make sure that gpu queue is empty
magma_device_sync();
magma_cgetmatrix_async( m-i*nb, nb, dA, cols, work, lda,
stream[0]);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), ldda,
dAT(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), ldda,
dAT(i, i-1), ldda,
c_one, dAT(i, i+1), ldda );
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lda, ipiv+i*nb, &iinfo);
}
if (*info == 0 && iinfo > 0)
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lda, dA, cols,
stream[0]);
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
magmablas_ctranspose( cols, nb, dA, cols, dAT(i,i), ldda );
// do the small non-parallel computations
if (s > (i+1)) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), ldda,
dAT(i+1, i ), ldda,
c_one, dAT(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), ldda,
dAT(i+1, i ), ldda,
c_one, dAT(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose( nb0, rows, dAT(s,s), ldda, dA, cols );
magma_cgetmatrix( rows, nb0, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo);
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb0, s*nb );
magma_csetmatrix( rows, nb0, work, lda, dA, cols );
magmablas_ctranspose( rows, nb0, dA, cols, dAT(s,s), ldda );
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s, s), ldda,
dAT(s, s)+nb0, ldda);
}
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_ctranspose_inplace( ldda, dAT, ldda );
magma_cgetmatrix( m, n, da, ldda, A, lda );
} else {
magmablas_ctranspose( n, m, dAT, ldda, da, maxm );
magma_cgetmatrix( m, n, da, maxm, A, lda );
magma_free( dAT );
}
magma_free( dA );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* magma_cgetrf */
extern "C" magma_int_t
magma_cgetrf_gpu(magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define dAT(i,j) (dAT + (i)*nb*lddat + (j)*nb)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, rows, cols, s, lddat, lddwork;
magmaFloatComplex *dAT, *dAP, *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
lddat = maxn;
lddwork = maxm;
dAT = dA;
if (MAGMA_SUCCESS != magma_cmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if ( m == n ) {
lddat = ldda;
magmablas_ctranspose_inplace( m, dAT, ldda );
}
else {
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ctranspose2( dAT, lddat, dA, ldda, m, n );
}
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
magma_free( dAP );
if ( ! (m == n))
magma_free( dAT );
*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;
for( i=0; i<s; i++ )
{
// download i-th panel
cols = maxm - i*nb;
//magmablas_ctranspose( dAP, cols, dAT(i,i), lddat, nb, cols );
magmablas_ctranspose2( dAP, cols, dAT(i,i), lddat, nb, m-i*nb );
// make sure that that the transpose has completed
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dAP, cols, work, lddwork,
stream[0]);
if ( i>0 ){
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), lddat,
dAT(i-1,i+1), lddat );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), lddat,
dAT(i, i-1), lddat,
c_one, dAT(i, i+1), lddat );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lddwork, ipiv+i*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dAP, maxm,
stream[0]);
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
//magmablas_ctranspose(dAT(i,i), lddat, dAP, maxm, cols, nb);
magmablas_ctranspose2(dAT(i,i), lddat, dAP, maxm, m-i*nb, nb);
// do the small non-parallel computations (next panel update)
if ( s > (i+1) ) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose2( dAP, maxm, dAT(s,s), lddat, nb0, rows);
magma_cgetmatrix( rows, nb0, dAP, maxm, work, lddwork );
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lddwork, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb0, s*nb );
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dAP, maxm );
magmablas_ctranspose2( dAT(s,s), lddat, dAP, maxm, rows, nb0);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
if ( m == n ) {
magmablas_ctranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ctranspose2( dA, ldda, dAT, lddat, n, m );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* End of MAGMA_CGETRF_GPU */
/**
Purpose
-------
CGETRF_m computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine. The matrix may exceed the GPU memory.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Note: The factorization of big panel is done calling multiple-gpu-interface.
Pivots are applied on GPU within the big panel.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_m(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magma_int_t *ipiv,
magma_int_t *info)
{
#define A(i,j) (A + (j)*lda + (i))
#define dAT(d,i,j) (dAT[d] + (i)*nb*ldn_local + (j)*nb)
#define dPT(d,i,j) (dPT[d] + (i)*nb*nb + (j)*nb*maxm)
magma_timer_t time=0, time_total=0, time_alloc=0, time_set=0, time_get=0, time_comp=0;
timer_start( time_total );
real_Double_t flops;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *dAT[MagmaMaxGPUs], *dA[MagmaMaxGPUs], *dPT[MagmaMaxGPUs];
magma_int_t iinfo = 0, nb, nbi, maxm, n_local[MagmaMaxGPUs], ldn_local;
magma_int_t N, M, NB, NBk, I, d, ngpu0 = ngpu;
magma_int_t ii, jj, h, offset, ib, rows;
magma_queue_t stream[MagmaMaxGPUs][2];
magma_event_t event[MagmaMaxGPUs][2];
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* initialize nb */
nb = magma_get_cgetrf_nb(m);
maxm = ((m + 31)/32)*32;
/* figure out NB */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
/* number of columns in the big panel */
h = 1+(2+ngpu0);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
//NB = 5*max(nb,32);
if ( ngpu0 > ceil((float)NB/nb) ) {
ngpu = (int)ceil((float)NB/nb);
h = 1+(2+ngpu);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
} else {
ngpu = ngpu0;
}
if ( ngpu*NB >= n ) {
#ifdef CHECK_CGETRF_OOC
printf( " * still fit in GPU memory.\n" );
#endif
NB = n;
} else {
#ifdef CHECK_CGETRF_OOC
printf( " * don't fit in GPU memory.\n" );
#endif
NB = ngpu*NB;
NB = max( nb, (NB / nb) * nb); /* making sure it's devisable by nb (x64) */
}
#ifdef CHECK_CGETRF_OOC
if ( NB != n ) printf( " * running in out-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n", n, NB, nb, (float)freeMem );
else printf( " * running in in-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n", n, NB, nb, (float)freeMem );
#endif
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code for scalar of one tile. */
lapackf77_cgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
/* allocate memory on GPU to store the big panel */
timer_start( time_alloc );
n_local[0] = (NB/nb)/ngpu;
if ( NB%(nb*ngpu) != 0 )
n_local[0]++;
n_local[0] *= nb;
ldn_local = ((n_local[0]+31)/32)*32;
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_cmalloc( &dA[d], (ldn_local+h*nb)*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dPT[d] = dA[d] + nb*maxm; /* for storing the previous panel from CPU */
dAT[d] = dA[d] + h*nb*maxm; /* for storing the big panel */
magma_queue_create( &stream[d][0] );
magma_queue_create( &stream[d][1] );
magma_event_create( &event[d][0] );
magma_event_create( &event[d][1] );
}
//magma_setdevice(0);
timer_stop( time_alloc );
for( I=0; I < n; I += NB ) {
M = m;
N = min( NB, n-I ); /* number of columns in this big panel */
//s = min( max(m-I,0), N )/nb; /* number of small block-columns in this big panel */
maxm = ((M + 31)/32)*32;
if ( ngpu0 > ceil((float)N/nb) ) {
ngpu = (int)ceil((float)N/nb);
} else {
ngpu = ngpu0;
}
for( d=0; d < ngpu; d++ ) {
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;
}
ldn_local = ((n_local[0]+31)/32)*32;
/* upload the next big panel into GPU, transpose (A->A'), and pivot it */
timer_start( time );
magmablas_csetmatrix_transpose_mgpu(ngpu, stream, A(0,I), lda,
dAT, ldn_local, dA, maxm, M, N, nb);
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
time_set += timer_stop( time );
timer_start( time );
/* == --------------------------------------------------------------- == */
/* == loop around the previous big-panels to update the new big-panel == */
for( offset = 0; offset < min(m,I); offset += NB ) {
NBk = min( m-offset, NB );
/* start sending the first tile from the previous big-panels to gpus */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
nbi = min( nb, NBk );
magma_csetmatrix_async( (M-offset), nbi,
A(offset,offset), lda,
dA[d], (maxm-offset), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][0] );
/* transpose */
magmablas_ctranspose( M-offset, nbi, dA[d], maxm-offset, dPT(d,0,0), nb );
}
/* applying the pivot from the previous big-panel */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magmablas_claswp_q( ldn_local, dAT(d,0,0), ldn_local, offset+1, offset+NBk, ipiv, 1, stream[d][1] );
}
/* == going through each block-column of previous big-panels == */
for( jj=0, ib=offset/nb; jj < NBk; jj += nb, ib++ ) {
ii = offset+jj;
rows = maxm - ii;
nbi = min( nb, NBk-jj );
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
/* wait for a block-column on GPU */
magma_queue_sync( stream[d][0] );
/* start sending next column */
if ( jj+nb < NBk ) {
magma_csetmatrix_async( (M-ii-nb), min(nb,NBk-jj-nb),
A(ii+nb,ii+nb), lda,
dA[d], (rows-nb), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][(1+jj/nb)%2] );
/* transpose next column */
magmablas_ctranspose( M-ii-nb, nb, dA[d], rows-nb, dPT(d,0,(1+jj/nb)%2), nb );
}
/* update with the block column */
magmablasSetKernelStream(stream[d][1]);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n_local[d], nbi, c_one, dPT(d,0,(jj/nb)%2), nb, dAT(d,ib,0), ldn_local );
if ( M > ii+nb ) {
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n_local[d], M-(ii+nb), nbi, c_neg_one, dAT(d,ib,0), ldn_local,
dPT(d,1,(jj/nb)%2), nb, c_one, dAT(d,ib+1,0), ldn_local );
}
magma_event_record( event[d][(jj/nb)%2], stream[d][1] );
} /* end of for each block-columns in a big-panel */
}
} /* end of for each previous big-panels */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
/* calling magma-gpu interface to panel-factorize the big panel */
if ( M > I ) {
magma_cgetrf2_mgpu(ngpu, M-I, N, nb, I, dAT, ldn_local, ipiv+I, dA, A(0,I), lda,
stream, &iinfo);
if ( iinfo < 0 ) {
*info = iinfo;
break;
} else if ( iinfo != 0 ) {
*info = iinfo + I * NB;
//break;
}
/* adjust pivots */
for( ii=I; ii < min(I+N,m); ii++ )
ipiv[ii] += I;
}
time_comp += timer_stop( time );
/* download the current big panel to CPU */
timer_start( time );
magmablas_cgetmatrix_transpose_mgpu(ngpu, stream, dAT, ldn_local, A(0,I), lda, dA, maxm, M, N, nb);
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
time_get += timer_stop( time );
} /* end of for */
timer_stop( time_total );
flops = FLOPS_CGETRF( m, n ) / 1e9;
timer_printf(" memory-allocation time: %e\n", time_alloc );
timer_printf(" NB=%d nb=%d\n", (int) NB, (int) nb );
timer_printf(" memcopy and transpose %e seconds\n", time_set );
timer_printf(" total time %e seconds\n", time_total );
timer_printf(" Performance %f GFlop/s, %f seconds without htod and dtoh\n", flops / (time_comp), time_comp );
timer_printf(" Performance %f GFlop/s, %f seconds with htod\n", flops / (time_comp + time_set), time_comp + time_set );
timer_printf(" Performance %f GFlop/s, %f seconds with dtoh\n", flops / (time_comp + time_get), time_comp + time_get );
timer_printf(" Performance %f GFlop/s, %f seconds without memory-allocation\n", flops / (time_total - time_alloc), time_total - time_alloc );
for( d=0; d < ngpu0; d++ ) {
magma_setdevice(d);
magma_free( dA[d] );
magma_event_destroy( event[d][0] );
magma_event_destroy( event[d][1] );
magma_queue_destroy( stream[d][0] );
magma_queue_destroy( stream[d][1] );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
if ( *info >= 0 )
magma_cgetrf_piv(m, n, NB, A, lda, ipiv, info);
return *info;
} /* magma_cgetrf_m */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cswap, cswapblk, claswp, claswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaFloatComplex *h_A1, *h_A2;
magmaFloatComplex *h_R1, *h_R2;
magmaFloatComplex_ptr d_A1, d_A2;
// row-major and column-major performance
real_Double_t row_perf0 = MAGMA_D_NAN, col_perf0 = MAGMA_D_NAN;
real_Double_t row_perf1 = MAGMA_D_NAN, col_perf1 = MAGMA_D_NAN;
real_Double_t row_perf2 = MAGMA_D_NAN, col_perf2 = MAGMA_D_NAN;
real_Double_t row_perf4 = MAGMA_D_NAN;
real_Double_t row_perf5 = MAGMA_D_NAN, col_perf5 = MAGMA_D_NAN;
real_Double_t row_perf6 = MAGMA_D_NAN, col_perf6 = MAGMA_D_NAN;
real_Double_t row_perf7 = MAGMA_D_NAN;
real_Double_t cpu_perf = MAGMA_D_NAN;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magmaInt_ptr d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" %8s cswap cswap cswapblk claswp claswp2 claswpx ccopymatrix CPU (all in )\n", g_platform_str );
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj/col-maj row-blk/col-blk claswp (GByte/s)\n");
printf("=========================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_cgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaFloatComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
// getrf always makes ipiv[j] >= j+1, where ipiv is one based and j is zero based
// some implementations (e.g., MacOS dlaswp) assume this
for( j=0; j < nb; j++ ) {
ipiv[j] = (rand() % (N-j)) + j + 1;
assert( ipiv[j] >= j+1 );
assert( ipiv[j] <= N );
}
/* =====================================================================
* cublas / clBLAS / Xeon Phi cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasCswap( opts.handle, N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1 );
#else
magma_cswap( N, d_A1, ldda*j, 1, d_A2, ldda*(ipiv[j]-1), 1, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
#ifdef HAVE_CUBLAS
cublasCswap( opts.handle, N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
#else
magma_cswap( N, d_A1, j, ldda, d_A2, ipiv[j]-1, ldda, opts.queue );
#endif
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswapblk, blocked version (2 matrices)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswp (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswp (1 matrix) - d_ipiv on GPU
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_claswp2( N, d_A1, ldda, 1, nb, d_ipiv, 1 );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* LAPACK-style claswpx (extended for row- and col-major) (1 matrix)
*/
#ifdef HAVE_CUBLAS
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
#endif
/* LAPACK swap on CPU for comparison */
time = magma_wtime();
lapackf77_claswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
#ifdef HAVE_CUBLAS
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
#endif
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_ccopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_ccopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf4, ((check & 0x040) != 0 ? '*' : ' '),
row_perf7, ((check & 0x080) != 0 ? '*' : ' '),
row_perf5, ((check & 0x100) != 0 ? '*' : ' '),
col_perf5, ((check & 0x200) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_cgetrf_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
magmaFloatComplex **d_lA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
=========
NUM_GPUS
(input) INTEGER
The number of GPUS to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define inAT(id,i,j) (d_lAT[(id)] + (i)*nb*lddat + (j)*nb)
magma_int_t nb, n_local[MagmaMaxGPUs];
magma_int_t maxm, mindim;
magma_int_t i, j, d, lddat, lddwork;
magmaFloatComplex *d_lAT[MagmaMaxGPUs];
magmaFloatComplex *d_panel[MagmaMaxGPUs], *work;
magma_queue_t streaml[MagmaMaxGPUs][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, d_lA[0], ldda, work, m );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, d_lA[0], ldda );
magma_free_cpu(work);
} else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if( num_gpus > ceil((float)n/nb) ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) num_gpus );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = ((((((n+nb-1)/nb)/num_gpus)*nb)+31)/32)*32;
lddat = (n+nb-1)/nb; /* number of block columns */
lddat = (lddat+num_gpus-1)/num_gpus; /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
lddat = ((lddat+31)/32)*32; /* make it a multiple of 32 */
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
/* local-n and local-ld */
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
/* workspaces */
if (MAGMA_SUCCESS != magma_cmalloc( &d_panel[i], (3+num_gpus)*nb*maxm )) {
for( j=0; j<=i; j++ ) {
magma_setdevice(j);
}
for( j=0; j<i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* local-matrix storage */
if (MAGMA_SUCCESS != magma_cmalloc( &d_lAT[i], lddat*maxm )) {
for( j=0; j<=i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
}
for( j=0; j<i; j++ ) {
magma_setdevice(j);
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* create the streams */
magma_queue_create( &streaml[i][0] );
magma_queue_create( &streaml[i][1] );
magmablasSetKernelStream(streaml[i][1]);
magmablas_ctranspose2( d_lAT[i], lddat, d_lA[i], ldda, m, n_local[i] );
}
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
cudaStreamSynchronize(streaml[i][0]);
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
/* cpu workspace */
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, lddwork*nb*num_gpus )) {
for(i=0; i<num_gpus; i++ ) {
magma_setdevice(i);
magma_free( d_panel[i] );
magma_free( d_lAT[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* calling multi-gpu interface with allocated workspaces and streams */
//magma_cgetrf1_mgpu( num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
// (magma_queue_t **)streaml, info );
magma_cgetrf2_mgpu(num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
streaml, info);
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
/* save on output */
magmablas_ctranspose2( d_lA[d], ldda, d_lAT[d], lddat, n_local[d], m );
magma_device_sync();
magma_free( d_lAT[d] );
magma_free( d_panel[d] );
magma_queue_destroy( streaml[d][0] );
magma_queue_destroy( streaml[d][1] );
magmablasSetKernelStream(NULL);
} /* end of for d=1,..,num_gpus */
magma_setdevice(0);
magma_free_pinned( work );
}
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cswap, cswapblk, cpermute, claswp, claswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaFloatComplex *h_A1, *h_A2;
magmaFloatComplex *d_A1, *d_A2;
magmaFloatComplex *h_R1, *h_R2;
// row-major and column-major performance
real_Double_t row_perf0, col_perf0;
real_Double_t row_perf1, col_perf1;
real_Double_t row_perf2, col_perf2;
real_Double_t row_perf3;
real_Double_t row_perf4;
real_Double_t row_perf5, col_perf5;
real_Double_t row_perf6, col_perf6;
real_Double_t row_perf7;
real_Double_t cpu_perf;
real_Double_t time, gbytes;
magma_int_t N, lda, ldda, nb, j;
magma_int_t ione = 1;
magma_int_t *ipiv, *ipiv2;
magma_int_t *d_ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_queue_t queue = 0;
printf(" cublasCswap cswap cswapblk claswp cpermute claswp2 claswpx ccopymatrix CPU (all in )\n");
printf(" N nb row-maj/col-maj row-maj/col-maj row-maj/col-maj row-maj row-maj row-maj row-maj/col-maj row-blk/col-blk claswp (GByte/s)\n");
printf("==================================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// For an N x N matrix, swap nb rows or nb columns using various methods.
// Each test is assigned one bit in the 'check' bitmask; bit=1 indicates failure.
// The variable 'shift' keeps track of which bit is for current test
int shift = 1;
int check = 0;
N = opts.nsize[itest];
lda = N;
ldda = ((N+31)/32)*32;
nb = (opts.nb > 0 ? opts.nb : magma_get_cgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(magmaFloatComplex) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R1, magmaFloatComplex, lda*N );
TESTING_MALLOC_PIN( h_R2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, nb );
TESTING_MALLOC_CPU( ipiv2, magma_int_t, nb );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, nb );
TESTING_MALLOC_DEV( d_A1, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_A2, magmaFloatComplex, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasCswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasCswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasCswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda);
}
}
time = magma_sync_wtime( queue ) - time;
col_perf0 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+ldda*j, 1, d_A2+ldda*(ipiv[j]-1), 1);
}
}
time = magma_sync_wtime( queue ) - time;
row_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
magmablas_cswap( N, d_A1+j, ldda, d_A2+ipiv[j]-1, ldda );
}
}
time = magma_sync_wtime( queue ) - time;
col_perf1 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaRowMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
row_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* Column Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_csetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_cswapblk( MagmaColMajor, N, d_A1, ldda, d_A2, ldda, 1, nb, ipiv, 1, 0);
time = magma_sync_wtime( queue ) - time;
col_perf2 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_cgetmatrix( N, N, d_A2, ldda, h_R2, lda );
check += (diff_matrix( N, N, h_A1, lda, h_R1, lda ) ||
diff_matrix( N, N, h_A2, lda, h_R2, lda ))*shift;
shift *= 2;
/* =====================================================================
* cpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // cpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_cpermute_long2( N, d_A1, ldda, ipiv2, nb, 0 );
time = magma_sync_wtime( queue ) - time;
row_perf3 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style claswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswp( N, d_A1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf4 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style claswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magma_setvector( nb, sizeof(magma_int_t), ipiv, 1, d_ipiv, 1 );
magmablas_claswp2( N, d_A1, ldda, 1, nb, d_ipiv, 1 );
time = magma_sync_wtime( queue ) - time;
row_perf7 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style claswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, ldda, 1, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
row_perf5 = gbytes / time;
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
blasf77_cswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* Col Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_csetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_claswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_claswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_cgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* Copy matrix.
*/
time = magma_sync_wtime( queue );
magma_ccopymatrix( N, nb, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
col_perf6 = 0.5 * gbytes / time;
time = magma_sync_wtime( queue );
magma_ccopymatrix( nb, N, d_A1, ldda, d_A2, ldda );
time = magma_sync_wtime( queue ) - time;
// copy reads 1 matrix and writes 1 matrix, so has half gbytes of swap
row_perf6 = 0.5 * gbytes / time;
printf("%5d %3d %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c/ %6.2f%c %6.2f%c %6.2f%c %6.2f%c %6.2f%c/ %6.2f%c %6.2f / %6.2f %6.2f %10s\n",
(int) N, (int) nb,
row_perf0, ((check & 0x001) != 0 ? '*' : ' '),
col_perf0, ((check & 0x002) != 0 ? '*' : ' '),
row_perf1, ((check & 0x004) != 0 ? '*' : ' '),
col_perf1, ((check & 0x008) != 0 ? '*' : ' '),
row_perf2, ((check & 0x010) != 0 ? '*' : ' '),
col_perf2, ((check & 0x020) != 0 ? '*' : ' '),
row_perf3, ((check & 0x040) != 0 ? '*' : ' '),
row_perf4, ((check & 0x080) != 0 ? '*' : ' '),
row_perf7, ((check & 0x100) != 0 ? '*' : ' '),
row_perf5, ((check & 0x200) != 0 ? '*' : ' '),
col_perf5, ((check & 0x400) != 0 ? '*' : ' '),
row_perf6,
col_perf6,
cpu_perf,
(check == 0 ? "ok" : "* failed") );
status += ! (check == 0);
TESTING_FREE_PIN( h_A1 );
TESTING_FREE_PIN( h_A2 );
TESTING_FREE_PIN( h_R1 );
TESTING_FREE_PIN( h_R2 );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( ipiv2 );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_A1 );
TESTING_FREE_DEV( d_A2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_cgetrf(magma_int_t m, magma_int_t n, cuFloatComplex *a, magma_int_t lda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. 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 = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
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,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define inAT(i,j) (dAT + (i)*nb*ldda + (j)*nb)
cuFloatComplex *dAT, *dA, *da, *work;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
nb = magma_get_cgetrf_nb(m);
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_cgetrf(&m, &n, a, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, maxdim;
magma_int_t i, rows, cols, s = min(m, n)/nb;
magma_int_t num_gpus = magma_num_gpus();
if ( num_gpus > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_int_t rval = magma_cgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
if( *info >= 0 ) magma_cgetrf_piv(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
ldda = maxn;
work = a;
if (maxdim*maxdim < 2*maxm*maxn)
{
if (MAGMA_SUCCESS != magma_cmalloc( &dA, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_int_t rval = magma_cgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
if( *info >= 0 ) magma_cgetrf_piv(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
ldda = maxdim;
magma_csetmatrix( m, n, a, lda, da, ldda );
dAT = da;
magmablas_cinplace_transpose( dAT, ldda, ldda );
}
else
{
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_int_t rval = magma_cgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
if( *info >= 0 ) magma_cgetrf_piv(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
magma_csetmatrix( m, n, a, lda, da, maxm );
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dA );
magma_int_t rval = magma_cgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
if( *info >= 0 ) magma_cgetrf_piv(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
magmablas_ctranspose2( dAT, ldda, da, maxm, m, n );
}
lapackf77_cgetrf( &m, &nb, work, &lda, ipiv, &iinfo);
for( i = 0; i < s; i++ )
{
// download i-th panel
cols = maxm - i*nb;
if (i>0){
magmablas_ctranspose( dA, cols, inAT(i,i), ldda, nb, cols );
magma_cgetmatrix( m-i*nb, nb, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, inAT(i-1,i-1), ldda,
inAT(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, inAT(i-1,i+1), ldda,
inAT(i, i-1), ldda,
c_one, inAT(i, i+1), ldda );
// do the cpu part
rows = m - i*nb;
lapackf77_cgetrf( &rows, &nb, work, &lda, ipiv+i*nb, &iinfo);
}
if (*info == 0 && iinfo > 0)
*info = iinfo + i*nb;
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb, i*nb );
// upload i-th panel
magma_csetmatrix( m-i*nb, nb, work, lda, dA, cols );
magmablas_ctranspose( inAT(i,i), ldda, dA, cols, cols, nb);
// do the small non-parallel computations
if (s > (i+1)){
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, inAT(i, i ), ldda,
inAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, inAT(i, i+1), ldda,
inAT(i+1, i ), ldda,
c_one, inAT(i+1, i+1), ldda );
}
else{
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, inAT(i, i ), ldda,
inAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, inAT(i, i+1), ldda,
inAT(i+1, i ), ldda,
c_one, inAT(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose2( dA, cols, inAT(s,s), ldda, nb0, rows);
magma_cgetmatrix( rows, nb0, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo);
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb0, s*nb );
magma_csetmatrix( rows, nb0, work, lda, dA, cols );
magmablas_ctranspose2( inAT(s,s), ldda, dA, cols, rows, nb0);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, inAT(s, s), ldda,
inAT(s, s)+nb0, ldda);
}
if (maxdim*maxdim< 2*maxm*maxn){
magmablas_cinplace_transpose( dAT, ldda, ldda );
magma_cgetmatrix( m, n, da, ldda, a, lda );
} else {
magmablas_ctranspose2( da, maxm, dAT, ldda, n, m );
magma_cgetmatrix( m, n, da, maxm, a, lda );
magma_free( dAT );
}
magma_free( dA );
}
return *info;
} /* magma_cgetrf */
/**
Purpose
-------
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
d_lA COMPLEX array of pointers on the GPU, dimension (ngpu).
On entry, the M-by-N matrix A distributed over GPUs
(d_lA[d] points to the local matrix on d-th GPU).
It uses 1D block column cyclic format with the block size of nb,
and each local matrix is stored by column.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array d_lA. LDDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_mgpu(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr d_lA[], magma_int_t ldda, magma_int_t *ipiv,
magma_int_t *info)
{
magma_int_t nb, n_local[MagmaMaxGPUs];
magma_int_t maxm;
magma_int_t i, j, d, lddat, lddwork;
magmaFloatComplex *d_lAT[MagmaMaxGPUs];
magmaFloatComplex *d_panel[MagmaMaxGPUs], *work;
magma_queue_t queues[MagmaMaxGPUs][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* create the queues */
for( d=0; d < ngpu; d++ ) {
magma_queue_create( d, &queues[d][0] );
magma_queue_create( d, &queues[d][1] );
}
/* Function Body */
nb = magma_get_cgetrf_nb( m, n );
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, d_lA[0], ldda, work, m, queues[0][0] );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, d_lA[0], ldda, queues[0][0] );
magma_free_cpu(work);
} else {
/* Use hybrid blocked code. */
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
maxm = magma_roundup( m, 32 );
if ( ngpu > ceil((float)n/nb) ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = magma_roundup( ((magma_ceildiv( n, nb )/ngpu)*nb), 32 );
lddat = magma_ceildiv( n, nb ); /* number of block columns */
lddat = magma_ceildiv( lddat, ngpu ); /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
lddat = magma_roundup( lddat, 32 ); /* make it a multiple of 32 */
for (i=0; i < ngpu; i++) {
magma_setdevice(i);
/* local-n and local-ld */
n_local[i] = ((n/nb)/ngpu)*nb;
if (i < (n/nb)%ngpu)
n_local[i] += nb;
else if (i == (n/nb)%ngpu)
n_local[i] += n%nb;
/* workspaces */
if (MAGMA_SUCCESS != magma_cmalloc( &d_panel[i], (3+ngpu)*nb*maxm )) {
for( j=0; j <= i; j++ ) {
magma_setdevice(j);
}
for( j=0; j < i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* local-matrix storage */
if (MAGMA_SUCCESS != magma_cmalloc( &d_lAT[i], lddat*maxm )) {
for( j=0; j <= i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
}
for( j=0; j < i; j++ ) {
magma_setdevice(j);
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ctranspose( m, n_local[i], d_lA[i], ldda, d_lAT[i], lddat, queues[i][1] );
}
for (i=0; i < ngpu; i++) {
magma_setdevice(i);
magma_queue_sync(queues[i][0]);
}
magma_setdevice(0);
/* cpu workspace */
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, lddwork*nb*ngpu )) {
for (i=0; i < ngpu; i++ ) {
magma_setdevice(i);
magma_free( d_panel[i] );
magma_free( d_lAT[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* calling multi-gpu interface with allocated workspaces and queues */
magma_cgetrf2_mgpu(ngpu, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
queues, info);
/* clean up */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
/* save on output */
magmablas_ctranspose( n_local[d], m, d_lAT[d], lddat, d_lA[d], ldda, queues[d][0] );
magma_queue_sync(queues[d][0]);
magma_queue_sync(queues[d][1]);
magma_free( d_lAT[d] );
magma_free( d_panel[d] );
} /* end of for d=1,..,ngpu */
magma_setdevice( orig_dev );
magma_free_pinned( work );
}
/* clean up */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_destroy( queues[d][0] );
magma_queue_destroy( queues[d][1] );
}
return *info;
}
extern "C" magma_int_t
magma_cgetrf_m(magma_int_t num_gpus0, magma_int_t m, magma_int_t n, magmaFloatComplex *a, magma_int_t lda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF_m computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine. The matrix may not fit entirely in the GPU memory.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Note: The factorization of big panel is done calling multiple-gpu-interface.
Pivots are applied on GPU within the big panel.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
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,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define A(i,j) (a + (j)*lda + (i))
#define inAT(d,i,j) (dAT[d] + (i)*nb*ldn_local + (j)*nb)
#define inPT(d,i,j) (dPT[d] + (i)*nb*nb + (j)*nb*maxm)
//#define PROFILE
#ifdef PROFILE
float flops, time_rmajor = 0, time_rmajor2 = 0, time_rmajor3 = 0, time_mem = 0;
magma_timestr_t start, start1, start2, end1, end, start0 = get_current_time();
#endif
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *dAT[MagmaMaxGPUs], *dA[MagmaMaxGPUs], *dPT[MagmaMaxGPUs];
magma_int_t iinfo = 0, nb, nbi, maxm, n_local[MagmaMaxGPUs], ldn_local;
magma_int_t N, M, NB, NBk, I, d, num_gpus;
magma_int_t ii, jj, h, offset, ib, rows, s;
magma_queue_t stream[MagmaMaxGPUs][2];
magma_event_t event[MagmaMaxGPUs][2];
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* initialize nb */
nb = magma_get_cgetrf_nb(m);
maxm = ((m + 31)/32)*32;
/* figure out NB */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
/* number of columns in the big panel */
h = 1+(2+num_gpus0);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
char * ngr_nb_char = getenv("MAGMA_NGR_NB");
if( ngr_nb_char != NULL ) NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
//NB = 5*max(nb,32);
if( num_gpus0 > ceil((float)NB/nb) ) {
num_gpus = (int)ceil((float)NB/nb);
h = 1+(2+num_gpus);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
} else {
num_gpus = num_gpus0;
}
if( num_gpus*NB >= n ) {
#ifdef CHECK_CGETRF_OOC
printf( " * still fit in GPU memory.\n" );
#endif
NB = n;
} else {
#ifdef CHECK_CGETRF_OOC
printf( " * don't fit in GPU memory.\n" );
#endif
NB = num_gpus*NB;
NB = max(nb,(NB / nb) * nb); /* making sure it's devisable by nb (x64) */
}
#ifdef CHECK_CGETRF_OOC
if( NB != n ) printf( " * running in out-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n",n,NB,nb,(float)freeMem );
else printf( " * running in in-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n",n,NB,nb,(float)freeMem );
#endif
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code for scalar of one tile. */
lapackf77_cgetrf(&m, &n, a, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
/* allocate memory on GPU to store the big panel */
#ifdef PROFILE
start = get_current_time();
#endif
n_local[0] = (NB/nb)/num_gpus;
if( NB%(nb*num_gpus) != 0 ) n_local[0] ++;
n_local[0] *= nb;
ldn_local = ((n_local[0]+31)/32)*32;
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_cmalloc( &dA[d], (ldn_local+h*nb)*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dPT[d] = dA[d] + nb*maxm; /* for storing the previous panel from CPU */
dAT[d] = dA[d] + h*nb*maxm; /* for storing the big panel */
magma_queue_create( &stream[d][0] );
magma_queue_create( &stream[d][1] );
magma_event_create( &event[d][0] );
magma_event_create( &event[d][1] );
}
//magma_setdevice(0);
#ifdef PROFILE
end = get_current_time();
printf( " memory-allocation time: %e\n",GetTimerValue(start, end)/1000.0 );
start = get_current_time();
#endif
for( I=0; I<n; I+=NB ) {
M = m;
N = min( NB, n-I ); /* number of columns in this big panel */
s = min(max(m-I,0),N)/nb; /* number of small block-columns in this big panel */
maxm = ((M + 31)/32)*32;
if( num_gpus0 > ceil((float)N/nb) ) {
num_gpus = (int)ceil((float)N/nb);
} else {
num_gpus = num_gpus0;
}
for( d=0; d<num_gpus; d++ ) {
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;
}
ldn_local = ((n_local[0]+31)/32)*32;
#ifdef PROFILE
start2 = get_current_time();
#endif
/* upload the next big panel into GPU, transpose (A->A'), and pivot it */
magmablas_csetmatrix_transpose_mgpu(num_gpus, stream, A(0,I), lda,
dAT, ldn_local, dA, maxm, M, N, nb);
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
#ifdef PROFILE
start1 = get_current_time();
#endif
/* == --------------------------------------------------------------- == */
/* == loop around the previous big-panels to update the new big-panel == */
for( offset = 0; offset<min(m,I); offset+=NB )
{
NBk = min( m-offset, NB );
/* start sending the first tile from the previous big-panels to gpus */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
nbi = min( nb, NBk );
magma_csetmatrix_async( (M-offset), nbi,
A(offset,offset), lda,
dA[d], (maxm-offset), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][0] );
/* transpose */
magmablas_ctranspose2( inPT(d,0,0), nb, dA[d], maxm-offset, M-offset, nbi);
}
/* applying the pivot from the previous big-panel */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][1]);
magmablas_cpermute_long3( inAT(d,0,0), ldn_local, ipiv, NBk, offset );
}
/* == going through each block-column of previous big-panels == */
for( jj=0, ib=offset/nb; jj<NBk; jj+=nb, ib++ )
{
ii = offset+jj;
rows = maxm - ii;
nbi = min( nb, NBk-jj );
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
/* wait for a block-column on GPU */
magma_queue_sync( stream[d][0] );
/* start sending next column */
if( jj+nb < NBk ) {
magma_csetmatrix_async( (M-ii-nb), min(nb,NBk-jj-nb),
A(ii+nb,ii+nb), lda,
dA[d], (rows-nb), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][(1+jj/nb)%2] );
/* transpose next column */
magmablas_ctranspose2( inPT(d,0,(1+jj/nb)%2), nb, dA[d], rows-nb, M-ii-nb, nb);
}
/* update with the block column */
magmablasSetKernelStream(stream[d][1]);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n_local[d], nbi, c_one, inPT(d,0,(jj/nb)%2), nb, inAT(d,ib,0), ldn_local );
if( M > ii+nb ) {
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n_local[d], M-(ii+nb), nbi, c_neg_one, inAT(d,ib,0), ldn_local,
inPT(d,1,(jj/nb)%2), nb, c_one, inAT(d,ib+1,0), ldn_local );
}
magma_event_record( event[d][(jj/nb)%2], stream[d][1] );
} /* end of for each block-columns in a big-panel */
}
} /* end of for each previous big-panels */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
/* calling magma-gpu interface to panel-factorize the big panel */
if( M > I ) {
//magma_cgetrf1_mgpu(num_gpus, M-I, N, nb, I, dAT, ldn_local, ipiv+I, dA, &a[I*lda], lda,
// (magma_queue_t **)stream, &iinfo);
magma_cgetrf2_mgpu(num_gpus, M-I, N, nb, I, dAT, ldn_local, ipiv+I, dA, A(0,I), lda,
stream, &iinfo);
if( iinfo < 0 ) {
*info = iinfo;
break;
} else if( iinfo != 0 ) {
*info = iinfo + I * NB;
//break;
}
/* adjust pivots */
for( ii=I; ii<min(I+N,m); ii++ )
ipiv[ii] += I;
}
#ifdef PROFILE
end1 = get_current_time();
time_rmajor += GetTimerValue(start1, end1);
time_rmajor3 += GetTimerValue(start2, end1);
time_mem += (GetTimerValue(start2, end1)-GetTimerValue(start1, end1))/1000.0;
#endif
/* download the current big panel to CPU */
magmablas_cgetmatrix_transpose_mgpu(num_gpus, stream, dAT, ldn_local, A(0,I), lda, dA, maxm, M, N, nb);
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
#ifdef PROFILE
end1 = get_current_time();
time_rmajor2 += GetTimerValue(start1, end1);
#endif
} /* end of for */
#ifdef PROFILE
end = get_current_time();
flops = FLOPS_CGETRF( m, n ) / 1000000;
printf(" NB=%d nb=%d\n",NB,nb);
printf(" memcopy and transpose %e seconds\n",time_mem );
printf(" total time %e seconds\n",GetTimerValue(start0,end)/1000.0);
printf(" Performance %f GFlop/s, %f seconds without htod and dtoh\n", flops / time_rmajor, time_rmajor /1000.0);
printf(" Performance %f GFlop/s, %f seconds with htod\n", flops / time_rmajor3, time_rmajor3/1000.0);
printf(" Performance %f GFlop/s, %f seconds with dtoh\n", flops / time_rmajor2, time_rmajor2/1000.0);
printf(" Performance %f GFlop/s, %f seconds without memory-allocation\n", flops / GetTimerValue(start, end), GetTimerValue(start,end)/1000.0);
#endif
for( d=0; d<num_gpus0; d++ ) {
magma_setdevice(d);
magma_free( dA[d] );
magma_event_destroy( event[d][0] );
magma_event_destroy( event[d][1] );
magma_queue_destroy( stream[d][0] );
magma_queue_destroy( stream[d][1] );
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
}
if( *info >= 0 ) magma_cgetrf_piv(m, n, NB, a, lda, ipiv, info);
return *info;
} /* magma_cgetrf_m */
/**
Purpose
-------
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@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, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *ipiv,
magma_int_t *info )
{
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, (dA_offset + (i_) + (j_)*ldda)
#define dAT(i_, j_) dAT, (dAT_offset + (i_)*lddat + (j_))
#define dAP(i_, j_) dAP, ( (i_) + (j_)*maxm)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#define dAT(i_, j_) (dAT + (i_)*lddat + (j_))
#define dAP(i_, j_) (dAP + (i_) + (j_)*maxm)
#endif
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, minmn;
magma_int_t i, j, jb, rows, lddat, ldwork;
magmaFloatComplex_ptr dAT=NULL, dAP=NULL;
magmaFloatComplex *work=NULL;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
minmn = min( m, n );
nb = magma_get_cgetrf_nb( m, n );
magma_queue_t queues[2] = { NULL };
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
if ( MAGMA_SUCCESS != magma_cmalloc_cpu( &work, m*n )) {
*info = MAGMA_ERR_HOST_ALLOC;
goto cleanup;
}
magma_cgetmatrix( m, n, dA(0,0), ldda, work, m, queues[0] );
lapackf77_cgetrf( &m, &n, work, &m, ipiv, info );
magma_csetmatrix( m, n, work, m, dA(0,0), ldda, queues[0] );
magma_free_cpu( work ); work=NULL;
}
else {
/* Use hybrid blocked code. */
maxm = magma_roundup( m, 32 );
maxn = magma_roundup( n, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto cleanup;
}
// square matrices can be done in place;
// rectangular requires copy to transpose
if ( m == n ) {
dAT = dA;
lddat = ldda;
magmablas_ctranspose_inplace( m, dAT(0,0), lddat, queues[0] );
}
else {
lddat = maxn; // N-by-M
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, lddat*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto cleanup;
}
magmablas_ctranspose( m, n, dA(0,0), ldda, dAT(0,0), lddat, queues[0] );
}
magma_queue_sync( queues[0] ); // finish transpose
ldwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, ldwork*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
goto cleanup;
}
for( j=0; j < minmn-nb; j += nb ) {
// get j-th panel from device
magmablas_ctranspose( nb, m-j, dAT(j,j), lddat, dAP(0,0), maxm, queues[1] );
magma_queue_sync( queues[1] ); // wait for transpose
magma_cgetmatrix_async( m-j, nb, dAP(0,0), maxm, work, ldwork, queues[0] );
if ( j > 0 ) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-(j+nb), nb,
c_one, dAT(j-nb, j-nb), lddat,
dAT(j-nb, j+nb), lddat, queues[1] );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+nb), m-j, nb,
c_neg_one, dAT(j-nb, j+nb), lddat,
dAT(j, j-nb), lddat,
c_one, dAT(j, j+nb), lddat, queues[1] );
}
// do the cpu part
rows = m - j;
magma_queue_sync( queues[0] ); // wait to get work
lapackf77_cgetrf( &rows, &nb, work, &ldwork, ipiv+j, &iinfo );
if ( *info == 0 && iinfo > 0 )
*info = iinfo + j;
// send j-th panel to device
magma_csetmatrix_async( m-j, nb, work, ldwork, dAP, maxm, queues[0] );
for( i=j; i < j + nb; ++i ) {
ipiv[i] += j;
}
magmablas_claswp( n, dAT(0,0), lddat, j + 1, j + nb, ipiv, 1, queues[1] );
magma_queue_sync( queues[0] ); // wait to set dAP
magmablas_ctranspose( m-j, nb, dAP(0,0), maxm, dAT(j,j), lddat, queues[1] );
// do the small non-parallel computations (next panel update)
if ( j + nb < minmn - nb ) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+nb), lddat, queues[1] );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(j+nb), nb,
c_neg_one, dAT(j, j+nb), lddat,
dAT(j+nb, j ), lddat,
c_one, dAT(j+nb, j+nb), lddat, queues[1] );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-(j+nb), nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+nb), lddat, queues[1] );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+nb), m-(j+nb), nb,
c_neg_one, dAT(j, j+nb), lddat,
dAT(j+nb, j ), lddat,
c_one, dAT(j+nb, j+nb), lddat, queues[1] );
}
}
jb = min( m-j, n-j );
if ( jb > 0 ) {
rows = m - j;
magmablas_ctranspose( jb, rows, dAT(j,j), lddat, dAP(0,0), maxm, queues[1] );
magma_cgetmatrix( rows, jb, dAP(0,0), maxm, work, ldwork, queues[1] );
// do the cpu part
lapackf77_cgetrf( &rows, &jb, work, &ldwork, ipiv+j, &iinfo );
if ( *info == 0 && iinfo > 0 )
*info = iinfo + j;
for( i=j; i < j + jb; ++i ) {
ipiv[i] += j;
}
magmablas_claswp( n, dAT(0,0), lddat, j + 1, j + jb, ipiv, 1, queues[1] );
// send j-th panel to device
magma_csetmatrix( rows, jb, work, ldwork, dAP(0,0), maxm, queues[1] );
magmablas_ctranspose( rows, jb, dAP(0,0), maxm, dAT(j,j), lddat, queues[1] );
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-j-jb, jb,
c_one, dAT(j,j), lddat,
dAT(j,j+jb), lddat, queues[1] );
}
// undo transpose
if ( m == n ) {
magmablas_ctranspose_inplace( m, dAT(0,0), lddat, queues[1] );
}
else {
magmablas_ctranspose( n, m, dAT(0,0), lddat, dA(0,0), ldda, queues[1] );
}
}
cleanup:
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dAP );
if (m != n) {
magma_free( dAT );
}
magma_free_pinned( work );
return *info;
} /* magma_cgetrf_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgetrf_mgpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float error;
magmaFloatComplex *h_A;
magmaFloatComplex *d_lA[ MagmaMaxGPUs ];
magma_int_t *ipiv;
magma_int_t M, N, n2, lda, ldda, n_local, ngpu;
magma_int_t info, min_mn, nb, ldn_local;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |Ax-b|/(N*|A|*|x|)\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |PA-LU|/(N*|A|)\n");
}
printf("=========================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_cgetrf_nb( M );
gflops = FLOPS_CGETRF( M, N ) / 1e9;
// ngpu must be at least the number of blocks
ngpu = min( opts.ngpu, int((N+nb-1)/nb) );
if ( ngpu < opts.ngpu ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
}
// Allocate host memory for the matrix
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
// Allocate device memory
for( int dev=0; dev < ngpu; dev++){
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
ldn_local = ((n_local+31)/32)*32; // TODO why?
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaFloatComplex, ldda*ldn_local );
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
init_matrix( M, N, h_A, lda );
cpu_time = magma_wtime();
lapackf77_cgetrf( &M, &N, h_A, &lda, ipiv, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
init_matrix( M, N, h_A, lda );
magma_csetmatrix_1D_col_bcyclic( M, N, h_A, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_cgetrf_mgpu( ngpu, M, N, d_lA, ldda, ipiv, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgetrf_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_A, lda, ngpu, nb );
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f)",
(int) M, (int) N, gpu_perf, gpu_time );
}
if ( opts.check == 2 ) {
error = get_residual( M, N, h_A, lda, ipiv );
printf(" %8.2e %s\n", error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else if ( opts.check ) {
error = get_LU_error( M, N, h_A, lda, ipiv );
printf(" %8.2e %s\n", error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
