extern "C" magma_int_t
magma_zgetrf_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
cuDoubleComplex **d_lA, 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
=======
ZGETRF 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_16 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)
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo, nb, n_local[MagmaMaxGPUs];
magma_int_t maxm, mindim;
magma_int_t i, j, d, rows, cols, s, lddat, lddwork;
magma_int_t id, i_local, i_local2, nb0, nb1;
cuDoubleComplex *d_lAT[MagmaMaxGPUs];
cuDoubleComplex *d_panel[MagmaMaxGPUs], *work;
cudaStream_t streaml[4][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_zgetrf_nb(m);
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_zmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_zgetmatrix( m, n, d_lA[0], ldda, work, m );
lapackf77_zgetrf(&m, &n, work, &m, ipiv, info);
magma_zsetmatrix( 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((double)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_zmalloc( &d_panel[i], 3*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_zmalloc( &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_ztranspose2( 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_zmalloc_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_zgetrf1_mgpu( num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
// (cudaStream_t **)streaml, info );
magma_zgetrf2_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_ztranspose2( 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;
/* End of MAGMA_ZGETRF_MGPU */
}
extern "C" magma_int_t
magma_zgetrf2_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n, magma_int_t nb, magma_int_t offset,
cuDoubleComplex **d_lAT, magma_int_t lddat, magma_int_t *ipiv,
cuDoubleComplex **d_lAP, cuDoubleComplex *w, magma_int_t ldw,
cudaStream_t streaml[][2], magma_int_t *info)
#endif
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2010
Purpose
=======
ZGETRF 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.
Use two buffer to send panels..
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_16 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
if INFO = -7, internal GPU 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)] + ((offset)+(i)*nb)*lddat + (j)*nb)
#define W(j) (w+((j)%num_gpus)*nb*ldw)
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t block_size = 32;
magma_int_t iinfo, n_local[4];
magma_int_t maxm, mindim;
magma_int_t i, ii, d, dd, rows, cols, s, ldpan[4];
magma_int_t id, i_local, i_local2, nb0, nb1;
cuDoubleComplex *d_panel[4], *panel_local[4];
//cudaStream_t streaml[4][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (num_gpus*lddat < max(1,n))
*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_zgetrf_nb(m);
if( num_gpus > ceil((double)n/nb) ) {
*info = -1;
return *info;
}
{
/* Use hybrid blocked code. */
maxm = ((m + block_size-1)/block_size)*block_size;
/* some initializations */
for(i=0; i<num_gpus; i++){
magmaSetDevice(i);
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 */
d_panel[i] = &(d_lAP[i][nb*maxm]); /* temporary panel storage */
/* create local streams */
//magma_queue_create(&streaml[i][0]);
//magma_queue_create(&streaml[i][1]);
}
trace_init( 1, num_gpus, 2, (CUstream_st**)streaml );
/* start sending the panel to cpu */
nb0 = min(mindim, nb);
magmaSetDevice(0);
magmablasSetKernelStream(streaml[0][1]);
trace_gpu_start( 0, 1, "comm", "get" );
if( nb0 == nb )
magmablas_ztranspose( d_lAP[0], maxm, inAT(0,0,0), lddat, nb0, maxm );
else
magmablas_ztranspose2( d_lAP[0], maxm, inAT(0,0,0), lddat, nb0, maxm );
magma_zgetmatrix_async( m, nb0,
d_lAP[0], maxm,
W(0), ldw, streaml[0][1] );
trace_gpu_end( 0, 1 );
/* ------------------------------------------------------------------------------------- */
#ifdef PROFILE
magma_timestr_t start_timer, end_timer;
start_timer = get_current_time();
#endif
s = mindim / nb;
for( i=0; i<s; i++ )
{
/* Set the GPU number that holds the current panel */
id = i%num_gpus;
magmaSetDevice(id);
/* Set the local index where the current panel is */
i_local = i/num_gpus;
cols = maxm - i*nb;
rows = m - i*nb;
/* synchrnoize i-th panel from id-th gpu into work */
magma_queue_sync( streaml[id][1] );
/* i-th panel factorization */
trace_cpu_start( 0, "getrf", "getrf" );
#ifdef PANEL_FACT_MC
cntxt->nb = 12;
magma_zgetrf_mc(cntxt, &rows, &nb, W(i), &ldw, ipiv+i*nb, &iinfo);
#else
lapackf77_zgetrf( &rows, &nb, W(i), &ldw, ipiv+i*nb, &iinfo);
#endif
if ( (*info == 0) && (iinfo > 0) ) {
*info = iinfo + i*nb;
//break;
}
trace_cpu_end( 0 );
/* start sending the panel to all the gpus */
d = (i+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
magmaSetDevice(d);
trace_gpu_start( 0, 1, "comm", "set" );
magma_zsetmatrix_async( rows, nb,
W(i), ldw,
d_lAP[d], cols, streaml[d][1] );
trace_gpu_end( 0, 1 );
d = (d+1)%num_gpus;
}
/* apply the pivoting */
d = (i+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][0]);
trace_gpu_start( d, 1, "pivot", "pivot" );
if( dd == 0 )
magmablas_zpermute_long2( lddat, inAT(d,0,0), lddat, ipiv, nb, i*nb );
else
magmablas_zpermute_long3( inAT(d,0,0), lddat, ipiv, nb, i*nb );
trace_gpu_end( d, 1 );
d = (d+1)%num_gpus;
}
/* update the trailing-matrix/look-ahead */
d = (i+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
magmaSetDevice(d);
/* storage for panel */
if( d == id ) {
/* the panel belond to this gpu */
panel_local[d] = inAT(d,i,i_local);
ldpan[d] = lddat;
/* next column */
i_local2 = i_local+1;
} else {
/* the panel belong to another gpu */
panel_local[d] = &d_panel[d][(i%2)*nb*maxm];
//panel_local[d] = d_panel[d];
ldpan[d] = nb;
/* next column */
i_local2 = i_local;
if( d < id ) i_local2 ++;
}
/* the size of the next column */
if ( s > (i+1) ) {
nb0 = nb;
} else {
nb0 = n_local[d]-nb*(s/num_gpus);
if( d < s%num_gpus ) nb0 -= nb;
}
if( d == (i+1)%num_gpus) {
/* owns the next column, look-ahead the column */
nb1 = nb0;
magmablasSetKernelStream(streaml[d][1]);
/* make sure all the pivoting has been applied */
magma_queue_sync(streaml[d][0]);
trace_gpu_start( d, 1, "gemm", "gemm" );
} else {
/* update the entire trailing matrix */
nb1 = n_local[d] - i_local2*nb;
magmablasSetKernelStream(streaml[d][0]);
/* synchronization to make sure panel arrived on gpu */
magma_queue_sync(streaml[d][1]);
trace_gpu_start( d, 0, "gemm", "gemm" );
}
magmablas_ztranspose(panel_local[d], ldpan[d], d_lAP[d], cols, cols, nb);
/* gpu updating the trailing matrix */
//magmablas_ztrsm(
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb1, nb, c_one,
panel_local[d], ldpan[d],
inAT(d, i, i_local2), lddat);
//cublasZgemm
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
nb1, m-(i+1)*nb, nb,
c_neg_one, inAT(d, i, i_local2), lddat,
&(panel_local[d][nb*ldpan[d]]), ldpan[d],
c_one, inAT(d, i+1, i_local2), lddat );
if( d == (i+1)%num_gpus )
{
/* Set the local index where the current panel is */
int loff = i+1;
int i_local = (i+1)/num_gpus;
int ldda = maxm - (i+1)*nb;
int cols = m - (i+1)*nb;
nb0 = min(nb, mindim - (i+1)*nb); /* size of the diagonal block */
trace_gpu_end( d, 1 );
if( nb0 > 0 ) {
/* transpose the panel for sending it to cpu */
trace_gpu_start( d, 1, "comm", "get" );
if( i+1 < s )
magmablas_ztranspose( d_lAP[d], ldda, inAT(d,loff,i_local), lddat, nb0, ldda );
else
magmablas_ztranspose2( d_lAP[d], ldda, inAT(d,loff,i_local), lddat, nb0, ldda );
/* send the panel to cpu */
magma_zgetmatrix_async( cols, nb0,
d_lAP[d], ldda,
W(i+1), ldw, streaml[d][1] );
trace_gpu_end( d, 1 );
}
} else {
trace_gpu_end( d, 0 );
}
d = (d+1)%num_gpus;
}
/* update the remaining matrix by gpu owning the next panel */
if( (i+1) < s ) {
int i_local = (i+1)/num_gpus;
int rows = m - (i+1)*nb;
d = (i+1)%num_gpus;
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][0]);
trace_gpu_start( d, 0, "gemm", "gemm" );
//magmablas_ztrsm
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n_local[d] - (i_local+1)*nb, nb,
c_one, panel_local[d], ldpan[d],
inAT(d,i,i_local+1), lddat );
//cublasZgemm
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n_local[d]-(i_local+1)*nb, rows, nb,
c_neg_one, inAT(d,i,i_local+1), lddat,
&(panel_local[d][nb*ldpan[d]]), ldpan[d],
c_one, inAT(d,i+1, i_local+1), lddat );
trace_gpu_end( d, 0 );
}
} /* end of for i=1..s */
/* ------------------------------------------------------------------------------ */
/* Set the GPU number that holds the last panel */
id = s%num_gpus;
/* Set the local index where the last panel is */
i_local = s/num_gpus;
/* size of the last diagonal-block */
nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
cols = maxm - s*nb;
if( nb0 > 0 ) {
magmaSetDevice(id);
/* wait for the last panel on cpu */
magma_queue_sync( streaml[id][1] );
/* factor on cpu */
lapackf77_zgetrf( &rows, &nb0, W(s), &ldw, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
/* send the factor to gpus */
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
i_local2 = i_local;
if( d < id ) i_local2 ++;
if( d == id || n_local[d] > i_local2*nb ) {
magma_zsetmatrix_async( rows, nb0,
W(s), ldw,
d_lAP[d], cols, streaml[d][1] );
}
}
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][0]);
if( d == 0 )
magmablas_zpermute_long2( lddat, inAT(d,0,0), lddat, ipiv, nb0, s*nb );
else
magmablas_zpermute_long3( inAT(d,0,0), lddat, ipiv, nb0, s*nb );
}
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][1]);
/* wait for the pivoting to be done */
magma_queue_sync( streaml[d][0] );
i_local2 = i_local;
if( d < id ) i_local2++;
if( d == id ) {
/* the panel belond to this gpu */
panel_local[d] = inAT(d,s,i_local);
/* next column */
nb1 = n_local[d] - i_local*nb-nb0;
magmablas_ztranspose2( panel_local[d], lddat, d_lAP[d], cols, rows, nb0);
if( nb1 > 0 )
//cublasZtrsm
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb1, nb0, c_one,
panel_local[d], lddat,
inAT(d,s,i_local)+nb0, lddat);
} else if( n_local[d] > i_local2*nb ) {
/* the panel belong to another gpu */
panel_local[d] = &d_panel[d][(s%2)*nb*maxm];
//panel_local[d] = d_panel[d];
/* next column */
nb1 = n_local[d] - i_local2*nb;
magmablas_ztranspose2( panel_local[d], nb, d_lAP[d], cols, rows, nb0);
//cublasZtrsm
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb1, nb0, c_one,
panel_local[d], nb,
inAT(d,s,i_local2), lddat);
}
}
} /* if( nb0 > 0 ) */
/* clean up */
trace_finalize( "zgetrf_mgpu.svg","trace.css" );
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
magma_queue_sync( streaml[d][0] );
magma_queue_sync( streaml[d][1] );
//magma_queue_destroy(streaml[d][0]);
//magma_queue_destroy(streaml[d][1]);
magmablasSetKernelStream(NULL);
}
magmaSetDevice(0);
#ifdef PROFILE
end_timer = get_current_time();
printf("\n Performance %f GFlop/s\n", (2./3.*n*n*n /1000000.) / GetTimerValue(start_timer, end_timer));
#endif
}
return *info;
/* End of MAGMA_ZGETRF2_MGPU */
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztranspose
Code is very similar to testing_zsymmetrize.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, gpu_perf2=0, gpu_time2=0, cpu_perf, cpu_time;
double error, error2, work[1];
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *h_A, *h_B, *h_R;
magmaDoubleComplex *d_A, *d_B;
magma_int_t M, N, size, lda, ldda, ldb, lddb;
magma_int_t ione = 1;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("Inplace transpose requires M==N.\n");
printf(" M N CPU GByte/s (sec) GPU GByte/s (sec) check Inplace GB/s (sec) check\n");
printf("====================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
lda = M;
ldda = ((M+31)/32)*32;
ldb = N;
lddb = ((N+31)/32)*32;
// load entire matrix, save entire matrix
gbytes = sizeof(magmaDoubleComplex) * 2.*M*N / 1e9;
// input is M x N
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*N );
// output is N x M
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*M );
TESTING_MALLOC( h_R, magmaDoubleComplex, ldb*M );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*M );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_Z_MAKE( i + j/10000., j );
}
}
for( int j = 0; j < M; ++j ) {
for( int i = 0; i < N; ++i ) {
h_B[i + j*ldb] = MAGMA_Z_MAKE( i + j/10000., j );
}
}
magma_zsetmatrix( N, M, h_B, ldb, d_B, lddb );
/* =====================================================================
Performs operation using naive out-of-place algorithm
(LAPACK doesn't implement transpose)
=================================================================== */
cpu_time = magma_wtime();
//for( int j = 1; j < N-1; ++j ) { // inset by 1 row & col
// for( int i = 1; i < M-1; ++i ) { // inset by 1 row & col
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_B[j + i*ldb] = h_A[i + j*lda];
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* ====================================================================
Performs operation using MAGMA, out-of-place
=================================================================== */
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
magma_zsetmatrix( N, M, h_B, ldb, d_B, lddb );
gpu_time = magma_sync_wtime( 0 );
//magmablas_ztranspose2( d_B+1+lddb, lddb, d_A+1+ldda, ldda, M-2, N-2 ); // inset by 1 row & col
magmablas_ztranspose2( d_B, lddb, d_A, ldda, M, N );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* ====================================================================
Performs operation using MAGMA, in-place
=================================================================== */
if ( M == N ) {
magma_zsetmatrix( M, N, h_A, lda, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 );
//magmablas_ztranspose_inplace( N-2, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_ztranspose_inplace( N, d_A, ldda );
gpu_time2 = magma_sync_wtime( 0 ) - gpu_time2;
gpu_perf2 = gbytes / gpu_time2;
}
/* =====================================================================
Check the result
=================================================================== */
size = ldb*M;
magma_zgetmatrix( N, M, d_B, lddb, h_R, ldb );
blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_zlange("f", &N, &M, h_R, &ldb, work );
if ( M == N ) {
magma_zgetmatrix( N, M, d_A, ldda, h_R, ldb );
blasf77_zaxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error2 = lapackf77_zlange("f", &N, &M, h_R, &ldb, work );
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %4s %7.2f (%7.2f) %4s\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
(error == 0. ? "ok" : "failed"),
gpu_perf2, gpu_time2,
(error2 == 0. ? "ok" : "failed") );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %4s --- ( --- )\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
(error == 0. ? "ok" : "failed") );
}
TESTING_FREE( h_A );
TESTING_FREE( h_B );
TESTING_FREE( h_R );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_zgetrf_gpu(magma_int_t m, magma_int_t n,
magmaDoubleComplex *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
=======
ZGETRF 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_16 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)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, rows, cols, s, lddat, lddwork;
magmaDoubleComplex *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_zgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_zmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_zgetmatrix( m, n, dA, ldda, work, m );
lapackf77_zgetrf(&m, &n, work, &m, ipiv, info);
magma_zsetmatrix( 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_zmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if ( m == n ) {
lddat = ldda;
magmablas_ztranspose_inplace( m, dAT, ldda );
}
else {
if (MAGMA_SUCCESS != magma_zmalloc( &dAT, maxm*maxn )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ztranspose2( dAT, lddat, dA, ldda, m, n );
}
if (MAGMA_SUCCESS != magma_zmalloc_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_ztranspose( dAP, cols, dAT(i,i), lddat, nb, cols );
magmablas_ztranspose2( dAP, cols, dAT(i,i), lddat, nb, m-i*nb );
// make sure that that the transpose has completed
magma_queue_sync( stream[1] );
magma_zgetmatrix_async( m-i*nb, nb, dAP, cols, work, lddwork,
stream[0]);
if ( i>0 ){
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), lddat,
dAT(i-1,i+1), lddat );
magma_zgemm( 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_zgetrf( &rows, &nb, work, &lddwork, ipiv+i*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_zsetmatrix_async( m-i*nb, nb, work, lddwork, dAP, maxm,
stream[0]);
magmablas_zpermute_long2( n, dAT, lddat, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
//magmablas_ztranspose(dAT(i,i), lddat, dAP, maxm, cols, nb);
magmablas_ztranspose2(dAT(i,i), lddat, dAP, maxm, m-i*nb, nb);
// do the small non-parallel computations (next panel update)
if ( s > (i+1) ) {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_zgemm( 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_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_zgemm( 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_ztranspose2( dAP, maxm, dAT(s,s), lddat, nb0, rows);
magma_zgetmatrix( rows, nb0, dAP, maxm, work, lddwork );
// do the cpu part
lapackf77_zgetrf( &rows, &nb0, work, &lddwork, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
magmablas_zpermute_long2( n, dAT, lddat, ipiv, nb0, s*nb );
// upload i-th panel
magma_zsetmatrix( rows, nb0, work, lddwork, dAP, maxm );
magmablas_ztranspose2( dAT(s,s), lddat, dAP, maxm, rows, nb0);
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
if ( m == n ) {
magmablas_ztranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ztranspose2( 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_ZGETRF_GPU */
extern "C" magma_int_t
magma_zgelqf( magma_int_t m, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda, magmaDoubleComplex *tau,
magmaDoubleComplex *work, magma_int_t lwork, magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZGELQF computes an LQ factorization of a COMPLEX_16 M-by-N matrix A:
A = L * Q.
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_16 array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and below the diagonal of the array
contain the m-by-min(m,n) lower trapezoidal matrix L (L is
lower triangular if m <= n); the elements above the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors (see Further Details).
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).
TAU (output) COMPLEX_16 array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= max(1,M).
For optimum performance LWORK >= M*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -10 internal GPU memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(k) . . . H(2) H(1), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:n) is stored on exit in A(i,i+1:n),
and tau in TAU(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a+(a_2)*(lda) + (a_1))
magmaDoubleComplex *dA, *dAT;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t maxm, maxn, maxdim, nb;
magma_int_t iinfo, ldda;
int lquery;
/* Function Body */
*info = 0;
nb = magma_get_zgelqf_nb(m);
work[0] = MAGMA_Z_MAKE( (double)(m*nb), 0 );
lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,m) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (min(m, n) == 0) {
work[0] = c_one;
return *info;
}
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
if (maxdim*maxdim < 2*maxm*maxn)
{
ldda = maxdim;
if (MAGMA_SUCCESS != magma_zmalloc( &dA, maxdim*maxdim )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_zsetmatrix( m, n, a, lda, dA, ldda );
dAT = dA;
magmablas_ztranspose_inplace( ldda, dAT, ldda );
}
else
{
ldda = maxn;
if (MAGMA_SUCCESS != magma_zmalloc( &dA, 2*maxn*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_zsetmatrix( m, n, a, lda, dA, maxm );
dAT = dA + maxn * maxm;
magmablas_ztranspose2( dAT, ldda, dA, maxm, m, n );
}
magma_zgeqrf2_gpu(n, m, dAT, ldda, tau, &iinfo);
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_ztranspose_inplace( ldda, dAT, ldda );
magma_zgetmatrix( m, n, dA, ldda, a, lda );
} else {
magmablas_ztranspose2( dA, maxm, dAT, ldda, n, m );
magma_zgetmatrix( m, n, dA, maxm, a, lda );
}
magma_free( dA );
return *info;
} /* magma_zgelqf */
