extern "C" magma_int_t
magma_dgetrf_gpu(magma_int_t m, magma_int_t n,
double *dA, magma_int_t ldda,
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
=======
DGETRF 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
=========
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) DOUBLE_PRECISION 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(i,j) (dAT + (i)*nb*lddat + (j)*nb)
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, rows, cols, s, lddat, lddwork;
double *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_dgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_dmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_dgetmatrix( m, n, dA, ldda, work, m );
lapackf77_dgetrf(&m, &n, work, &m, ipiv, info);
magma_dsetmatrix( 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_dmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if ((m == n) && (m % 32 == 0) && (ldda%32 == 0)){
lddat = ldda;
magmablas_dinplace_transpose( dAT, ldda, m);
}
else {
if (MAGMA_SUCCESS != magma_dmalloc( &dAT, maxm*maxn )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_dtranspose2( dAT, lddat, dA, ldda, m, n );
}
if (MAGMA_SUCCESS != magma_dmalloc_pinned( &work, maxm*nb )) {
magma_free( dAP );
if (! ((m == n) && (m % 32 == 0) && (ldda%32 == 0)) )
magma_free( dAT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
for( i=0; i<s; i++ )
{
// download i-th panel
cols = maxm - i*nb;
magmablas_dtranspose( dAP, cols, inAT(i,i), lddat, nb, cols );
magma_dgetmatrix( m-i*nb, nb, dAP, cols, work, lddwork );
// make sure that gpu queue is empty
magma_device_sync();
if ( i>0 ){
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, inAT(i-1,i-1), lddat,
inAT(i-1,i+1), lddat );
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, inAT(i-1,i+1), lddat,
inAT(i, i-1), lddat,
c_one, inAT(i, i+1), lddat );
}
// do the cpu part
rows = m - i*nb;
lapackf77_dgetrf( &rows, &nb, work, &lddwork, ipiv+i*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
magmablas_dpermute_long2( n, dAT, lddat, ipiv, nb, i*nb );
// upload i-th panel
magma_dsetmatrix( m-i*nb, nb, work, lddwork, dAP, maxm );
magmablas_dtranspose(inAT(i,i), lddat, dAP, maxm, cols, nb);
// do the small non-parallel computations
if ( s > (i+1) ) {
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, inAT(i, i ), lddat,
inAT(i, i+1), lddat);
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, inAT(i, i+1), lddat,
inAT(i+1, i ), lddat,
c_one, inAT(i+1, i+1), lddat );
}
else {
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, inAT(i, i ), lddat,
inAT(i, i+1), lddat);
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, inAT(i, i+1), lddat,
inAT(i+1, i ), lddat,
c_one, inAT(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_dtranspose2( dAP, maxm, inAT(s,s), lddat, nb0, rows);
magma_dgetmatrix( rows, nb0, dAP, maxm, work, lddwork );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_dgetrf( &rows, &nb0, work, &lddwork, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
magmablas_dpermute_long2( n, dAT, lddat, ipiv, nb0, s*nb );
// upload i-th panel
magma_dsetmatrix( rows, nb0, work, lddwork, dAP, maxm );
magmablas_dtranspose2( inAT(s,s), lddat, dAP, maxm, rows, nb0);
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, inAT(s,s), lddat,
inAT(s,s)+nb0, lddat);
if ((m == n) && (m % 32 == 0) && (ldda%32 == 0)){
magmablas_dinplace_transpose( dAT, lddat, m );
}
else {
magmablas_dtranspose2( dA, ldda, dAT, lddat, n, m );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
}
return *info;
/* End of MAGMA_DGETRF_GPU */
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dswap, dswapblk, dpermute, dlaswp, dlaswpx
*/
int main( int argc, char** argv)
{
TESTING_INIT();
double *h_A1, *h_A2;
double *d_A1, *d_A2;
double *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(" cublasDswap dswap dswapblk dlaswp dpermute dlaswp2 dlaswpx dcopymatrix 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 dlaswp (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_dgetrf_nb( N ));
nb = min( N, nb );
// each swap does 2N loads and 2N stores, for nb swaps
gbytes = sizeof(double) * 4.*N*nb / 1e9;
TESTING_MALLOC_PIN( h_A1, double, lda*N );
TESTING_MALLOC_PIN( h_A2, double, lda*N );
TESTING_MALLOC_PIN( h_R1, double, lda*N );
TESTING_MALLOC_PIN( h_R2, double, 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, double, ldda*N );
TESTING_MALLOC_DEV( d_A2, double, ldda*N );
for( j=0; j < nb; j++ ) {
ipiv[j] = (magma_int_t) ((rand()*1.*N) / (RAND_MAX * 1.)) + 1;
}
/* =====================================================================
* cublasDswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasDswap( 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_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
for( j=0; j < nb; j++) {
if ( j != (ipiv[j]-1)) {
cublasDswap( 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_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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;
/* =====================================================================
* dswap, row-by-row (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( 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_dswap( 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_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( 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_dswap( 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_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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;
/* =====================================================================
* dswapblk, blocked version (2 matrices)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
init_matrix( N, N, h_A2, lda, 100 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_dswapblk( 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_dswap( &N, h_A1+lda*j, &ione, h_A2+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
magma_dsetmatrix( N, N, h_A2, lda, d_A2, ldda );
time = magma_sync_wtime( queue );
magmablas_dswapblk( 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_dswap( &N, h_A1+j, &lda, h_A2+(ipiv[j]-1), &lda);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
magma_dgetmatrix( 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;
/* =====================================================================
* dpermute_long (1 matrix)
*/
/* Row Major */
memcpy( ipiv2, ipiv, nb*sizeof(magma_int_t) ); // dpermute updates ipiv2
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dpermute_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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswp (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswp( 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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswp (1 matrix) - d_ipiv on GPU
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( 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_dlaswp2( 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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( N, N, d_A1, ldda, h_R1, lda );
check += diff_matrix( N, N, h_A1, lda, h_R1, lda )*shift;
shift *= 2;
/* =====================================================================
* LAPACK-style dlaswpx (extended for row- and col-major) (1 matrix)
*/
/* Row Major */
init_matrix( N, N, h_A1, lda, 0 );
magma_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswpx( 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_dswap( &N, h_A1+lda*j, &ione, h_A1+lda*(ipiv[j]-1), &ione);
}
}
magma_dgetmatrix( 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_dsetmatrix( N, N, h_A1, lda, d_A1, ldda );
time = magma_sync_wtime( queue );
magmablas_dlaswpx( N, d_A1, 1, ldda, 1, nb, ipiv, 1);
time = magma_sync_wtime( queue ) - time;
col_perf5 = gbytes / time;
time = magma_wtime();
lapackf77_dlaswp( &N, h_A1, &lda, &ione, &nb, ipiv, &ione);
time = magma_wtime() - time;
cpu_perf = gbytes / time;
magma_dgetmatrix( 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_dcopymatrix( 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_dcopymatrix( 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_dgetrf(magma_int_t m, magma_int_t n, double *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
=======
DGETRF 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
=========
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) DOUBLE_PRECISION 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 dAT(i,j) (dAT + (i)*nb*ldda + (j)*nb)
double *dAT, *dA, *da, *work;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_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_dgetrf_nb(m);
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_dgetrf(&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_dgetrf_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(double);
int h = 1+(2+num_gpus), num_gpus2 = num_gpus;
int 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) ) );
if( num_gpus > ceil((double)NB/nb) ) {
num_gpus2 = (int)ceil((double)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_dgetrf_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_dmalloc( &dA, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_dgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
ldda = maxdim;
magma_dsetmatrix( m, n, a, lda, da, ldda );
dAT = da;
magmablas_dtranspose_inplace( ldda, dAT, ldda );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
if (MAGMA_SUCCESS != magma_dmalloc( &dA, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_dgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
magma_dsetmatrix( m, n, a, lda, da, maxm );
if (MAGMA_SUCCESS != magma_dmalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dA );
magma_dgetrf_m(num_gpus, m, n, a, lda, ipiv, info);
return *info;
}
magmablas_dtranspose2( dAT, ldda, da, maxm, m, n );
}
lapackf77_dgetrf( &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_dtranspose( dA, cols, dAT(i,i), ldda, nb, cols );
// make sure that gpu queue is empty
magma_device_sync();
magma_dgetmatrix_async( m-i*nb, nb, dA, cols, work, lda,
stream[0]);
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), ldda,
dAT(i-1,i+1), ldda );
magma_dgemm( 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_dgetrf( &rows, &nb, work, &lda, ipiv+i*nb, &iinfo);
}
if (*info == 0 && iinfo > 0)
*info = iinfo + i*nb;
// upload i-th panel
magma_dsetmatrix_async( m-i*nb, nb, work, lda, dA, cols,
stream[0]);
magmablas_dpermute_long2( ldda, dAT, ldda, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
magmablas_dtranspose( dAT(i,i), ldda, dA, cols, cols, nb);
// do the small non-parallel computations
if (s > (i+1)){
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_dgemm( 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_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_dgemm( 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_dtranspose2( dA, cols, dAT(s,s), ldda, nb0, rows);
magma_dgetmatrix( rows, nb0, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_dgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo);
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
magmablas_dpermute_long2( ldda, dAT, ldda, ipiv, nb0, s*nb );
magma_dsetmatrix( rows, nb0, work, lda, dA, cols );
magmablas_dtranspose2( dAT(s,s), ldda, dA, cols, rows, nb0);
magma_dtrsm( 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_dtranspose_inplace( ldda, dAT, ldda );
magma_dgetmatrix( m, n, da, ldda, a, lda );
} else {
magmablas_dtranspose2( da, maxm, dAT, ldda, n, m );
magma_dgetmatrix( 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_dgetrf */
