extern "C" magma_int_t
magma_dgeqrs_gpu(
magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDouble_ptr dA, size_t dA_offset, magma_int_t ldda,
double *tau, magmaDouble_ptr dT, size_t dT_offset,
magmaDouble_ptr dB, size_t dB_offset, magma_int_t lddb,
double *hwork, magma_int_t lwork,
magma_queue_t queue,
magma_int_t *info)
{
/* -- clMagma (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by DGEQRF_GPU.
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. M >= N >= 0.
NRHS (input) INTEGER
The number of columns of the matrix C. NRHS >= 0.
A (input) DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
DGEQRF_GPU in the first n columns of its array argument A.
LDDA (input) INTEGER
The leading dimension of the array A, LDDA >= M.
TAU (input) DOUBLE_PRECISION array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_DGEQRF_GPU.
DB (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
DT (input) DOUBLE_PRECISION array that is the output (the 6th argument)
of magma_dgeqrf_gpu of size
2*MIN(M, N)*NB + ((N+31)/32*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
((N+31)/32*32 )* MAX(NB, NRHS).
LDDB (input) INTEGER
The leading dimension of the array DB. LDDB >= M.
HWORK (workspace/output) DOUBLE_PRECISION array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK,
LWORK >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by magma_get_dgeqrf_nb( M ).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the WORK array.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
#define a_ref(a_1,a_2) dA, (dA_offset + (a_1) + (a_2)*(ldda))
#define d_ref(a_1) dT, (dT_offset + (lddwork+(a_1))*nb)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magmaDouble_ptr dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_dgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_D_MAKE( (double)lwkopt, 0. );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -8;
else if (lwork < lwkopt && ! lquery)
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
hwork[0] = c_one;
return *info;
}
/* B := Q' * B */
magma_dormqr_gpu( MagmaLeft, MagmaConjTrans,
m, nrhs, n,
a_ref(0,0), ldda, tau,
dB, dB_offset, lddb, hwork, lwork, dT, dT_offset, nb, queue, info );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= k;
int ldtwork;
size_t dwork_offset = 0;
if (nb < k) {
dwork = dT;
dwork_offset = dT_offset+2*lddwork*nb;
}
else {
ldtwork = ( 2*k + ((n+31)/32)*32 )*nb;
magma_dmalloc( &dwork, ldtwork );
}
// To do: Why did we have this line originally; seems to be a bug (Stan)?
//dwork = dT;
i = (k-1)/nb * nb;
ib = n-i;
rows = m-i;
// TODO: this assumes that, on exit from magma_dormqr_gpu, hwork contains
// the last block of A and B (i.e., C in dormqr). This should be fixed.
// Seems this data should already be on the GPU, so could switch to
// magma_dtrsm and drop the dsetmatrix.
if ( nrhs == 1 ) {
blasf77_dtrsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_dtrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_dsetmatrix( ib, nrhs, hwork+rows*ib, rows, dwork, dwork_offset+i, lddwork, queue );
// update c
if (nrhs == 1)
magma_dgemv( MagmaNoTrans, i, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, 1,
c_one, dB, dB_offset, 1, queue );
else
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset + i, lddwork,
c_one, dB, dB_offset, lddb, queue );
int start = i-nb;
if (nb < k) {
for (i = start; i >=0; i -= nb) {
ib = min(k-i, nb);
rows = m -i;
if (i + ib < n) {
if (nrhs == 1) {
magma_dgemv( MagmaNoTrans, ib, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, 1,
c_zero, dwork, dwork_offset+i, 1, queue );
magma_dgemv( MagmaNoTrans, i, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, 1,
c_one, dB, dB_offset, 1, queue );
}
else {
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
ib, nrhs, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, lddb,
c_zero, dwork, dwork_offset+i, lddwork, queue );
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, lddwork,
c_one, dB, dB_offset, lddb, queue );
}
}
}
}
magma_dcopymatrix( (n), nrhs,
dwork, dwork_offset, lddwork,
dB, dB_offset, lddb, queue );
if (nb >= k)
magma_free(dwork);
magma_queue_sync( queue );
return *info;
}
/**
Purpose
-------
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by DGEQRF_GPU.
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. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in]
dA DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
DGEQRF_GPU in the first n columns of its array argument A.
@param[in]
ldda INTEGER
The leading dimension of the array A, LDDA >= M.
@param[in]
tau DOUBLE_PRECISION array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_DGEQRF_GPU.
@param[in,out]
dB DOUBLE_PRECISION array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in]
dT DOUBLE_PRECISION array that is the output (the 6th argument)
of magma_dgeqrf_gpu of size
2*MIN(M, N)*NB + ((N+31)/32*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
((N+31)/32*32 )* MAX(NB, NRHS).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) DOUBLE_PRECISION array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK,
LWORK >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by magma_get_dgeqrf_nb( M ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the WORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_dgels_comp
********************************************************************/
extern "C" magma_int_t
magma_dgeqrs_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
double *dA, magma_int_t ldda,
double *tau, double *dT,
double *dB, magma_int_t lddb,
double *hwork, magma_int_t lwork,
magma_int_t *info)
{
#define dA(a_1,a_2) (dA + (a_2)*(ldda) + (a_1))
#define dT(a_1) (dT + (lddwork+(a_1))*nb)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_dgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_D_MAKE( (double)lwkopt, 0. );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -9;
else if (lwork < lwkopt && ! lquery)
*info = -11;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
hwork[0] = c_one;
return *info;
}
/* B := Q' * B */
magma_dormqr_gpu( MagmaLeft, MagmaTrans,
m, nrhs, n,
dA(0,0), ldda, tau,
dB, lddb, hwork, lwork, dT, nb, info );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= k;
if (nb < k)
dwork = dT+2*lddwork*nb;
else
dwork = dT;
// To do: Why did we have this line originally; seems to be a bug (Stan)?
// dwork = dT;
i = (k-1)/nb * nb;
ib = n-i;
rows = m-i;
// TODO: this assumes that, on exit from magma_dormqr_gpu, hwork contains
// the last block of A and B (i.e., C in dormqr). This should be fixed.
// Seems this data should already be on the GPU, so could switch to
// magma_dtrsm and drop the dsetmatrix.
if ( nrhs == 1 ) {
blasf77_dtrsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_dtrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_dsetmatrix( ib, nrhs, hwork+rows*ib, rows, dwork+i, lddwork );
// update c
if (nrhs == 1)
magma_dgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1);
else
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb);
int start = i-nb;
if (nb < k) {
for (i = start; i >= 0; i -= nb) {
ib = min(k-i, nb);
rows = m -i;
if (i + ib < n) {
if (nrhs == 1) {
magma_dgemv( MagmaNoTrans, ib, ib,
c_one, dT(i), ib,
dB+i, 1,
c_zero, dwork+i, 1);
magma_dgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1);
}
else {
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
ib, nrhs, ib,
c_one, dT(i), ib,
dB+i, lddb,
c_zero, dwork+i, lddwork);
magma_dgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb);
}
}
}
}
magma_dcopymatrix( (n), nrhs,
dwork, lddwork,
dB, lddb );
return *info;
}
/***************************************************************************//**
Purpose
-------
DGELS solves the overdetermined, least squares problem
min || A*X - C ||
using the QR factorization A.
The underdetermined problem (m < n) is not currently handled.
Arguments
---------
@param[in]
trans magma_trans_t
- = MagmaNoTrans: the linear system involves A.
Only TRANS=MagmaNoTrans is currently handled.
@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. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, A is overwritten by details of its QR
factorization as returned by DGEQRF.
@param[in]
lda INTEGER
The leading dimension of the array A, LDA >= M.
@param[in,out]
B DOUBLE PRECISION array, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= M.
@param[out]
hwork (workspace) DOUBLE PRECISION array, dimension MAX(1,LWORK).
On exit, if INFO = 0, HWORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array HWORK,
LWORK >= max( N*NB, 2*NB*NB ),
where NB is the blocksize given by magma_get_dgeqrf_nb( M, N ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the HWORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_gels
*******************************************************************************/
extern "C" magma_int_t
magma_dgels(
magma_trans_t trans, magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDouble_ptr A, magma_int_t lda,
magmaDouble_ptr B, magma_int_t ldb,
double *hwork, magma_int_t lwork,
magma_int_t *info)
{
/* Constants */
const double c_one = MAGMA_D_ONE;
/* Local variables */
double *tau;
magma_int_t min_mn;
magma_int_t nb = magma_get_dgeqrf_nb( m, n );
magma_int_t lwkopt = max( n*nb, 2*nb*nb ); // (m - n + nb)*(nrhs + nb) + nrhs*nb;
bool lquery = (lwork == -1);
hwork[0] = magma_dmake_lwork( lwkopt );
*info = 0;
/* For now, N is the only case working */
if ( trans != MagmaNoTrans )
*info = -1;
else if (m < 0)
*info = -2;
else if (n < 0 || m < n) /* LQ is not handle for now */
*info = -3;
else if (nrhs < 0)
*info = -4;
else if (lda < max(1,m))
*info = -6;
else if (ldb < max(1,m))
*info = -8;
else if (lwork < lwkopt && ! lquery)
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
min_mn = min(m,n);
if (min_mn == 0) {
hwork[0] = c_one;
return *info;
}
magma_dmalloc_cpu( &tau, min_mn );
if ( tau == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_dgeqrf( m, n, A, lda, tau, hwork, lwork, info );
if ( *info == 0 ) {
// B := Q' * B
lapackf77_dormqr( MagmaLeftStr, MagmaTransStr, &m, &nrhs, &n,
A, &lda, tau, B, &ldb, hwork, &lwork, info );
// Solve R*X = B(1:n,:)
blasf77_dtrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&n, &nrhs, &c_one, A, &lda, B, &ldb );
}
magma_free_cpu( tau );
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dtrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double magma_error, cublas_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *piv;
magma_err_t err;
double *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2, *LU, *LUT;
double *d_A, *d_B;
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double alpha = MAGMA_D_MAKE( 0.29, -0.86 );
magma_opts opts;
parse_opts( argc, argv, &opts );
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"side = %c, uplo = %c, transA = %c, diag = %c \n", opts.side, opts.uplo, opts.transA, opts.diag );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\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];
gflops = FLOPS_DTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC( h_A, double, lda*Ak );
TESTING_MALLOC( LU, double, lda*Ak );
TESTING_MALLOC( LUT, double, lda*Ak );
TESTING_MALLOC( h_B, double, ldb*N );
TESTING_MALLOC( h_B1, double, ldb*N );
TESTING_MALLOC( h_X1, double, ldb*N );
TESTING_MALLOC( h_X2, double, ldb*N );
TESTING_MALLOC( h_Bcublas, double, ldb*N );
TESTING_MALLOC( h_Bmagma, double, ldb*N );
TESTING_DEVALLOC( d_A, double, ldda*Ak );
TESTING_DEVALLOC( d_B, double, lddb*N );
/* Initialize the matrices */
lapackf77_dlarnv( &ione, ISEED, &sizeA, LU );
err = magma_malloc_cpu( (void**) &piv, Ak*sizeof(magma_int_t) ); assert( err == 0 );
lapackf77_dgetrf( &Ak, &Ak, LU, &lda, piv, &info );
int i, j;
for(i=0;i<Ak;i++){
for(j=0;j<Ak;j++){
LUT[j+i*lda] = LU[i+j*lda];
}
}
lapackf77_dlacpy(MagmaUpperStr, &Ak, &Ak, LUT, &lda, LU, &lda);
if(opts.uplo == MagmaLower){
lapackf77_dlacpy(MagmaLowerStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}else{
lapackf77_dlacpy(MagmaUpperStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}
lapackf77_dlarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(double));
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_dsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_dtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_dgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_dsetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasDtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_dgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_dtrsm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(double));
double alpha2 = MAGMA_D_DIV( c_one, alpha );
blasf77_dtrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_daxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
double norm1 = lapackf77_dlange( "M", &M, &N, h_X1, &ldb, work );
double normx = lapackf77_dlange( "M", &M, &N, h_Bmagma, &ldb, work );
double normA = lapackf77_dlange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(double));
blasf77_dtrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_daxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_dlange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_dlange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_dlange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error );
}
TESTING_FREE( h_A );
TESTING_FREE( LU );
TESTING_FREE( LUT );
TESTING_FREE( h_B );
TESTING_FREE( h_Bcublas );
TESTING_FREE( h_Bmagma );
TESTING_FREE( h_B1 );
TESTING_FREE( h_X1 );
TESTING_FREE( h_X2 );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/**
Purpose
-------
DGETRF_NOPIV computes an LU factorization of a general M-by-N
matrix A without pivoting.
The factorization has the form
A = L * U
where 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]
A 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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 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_dgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_dgetrf_nopiv(
magma_int_t m, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t min_mn, i__3, i__4;
magma_int_t j, jb, nb, iinfo;
A -= 1 + lda;
/* Function Body */
*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;
}
/* Determine the block size for this environment. */
nb = 128;
min_mn = min(m,n);
if (nb <= 1 || nb >= min_mn) {
/* Use unblocked code. */
magma_dgetf2_nopiv( m, n, A(1,1), lda, info );
}
else {
/* Use blocked code. */
for (j = 1; j <= min_mn; j += nb) {
jb = min( min_mn - j + 1, nb );
/* Factor diagonal and subdiagonal blocks and test for exact
singularity. */
i__3 = m - j + 1;
//magma_dgetf2_nopiv( i__3, jb, A(j,j), lda, &iinfo );
i__3 -= jb;
magma_dgetf2_nopiv( jb, jb, A(j,j), lda, &iinfo );
blasf77_dtrsm( "R", "U", "N", "N", &i__3, &jb, &c_one,
A(j,j), &lda,
A(j+jb,j), &lda );
/* Adjust INFO */
if (*info == 0 && iinfo > 0)
*info = iinfo + j - 1;
if (j + jb <= n) {
/* Compute block row of U. */
i__3 = n - j - jb + 1;
blasf77_dtrsm( "Left", "Lower", "No transpose", "Unit",
&jb, &i__3, &c_one,
A(j,j), &lda,
A(j,j+jb), &lda );
if (j + jb <= m) {
/* Update trailing submatrix. */
i__3 = m - j - jb + 1;
i__4 = n - j - jb + 1;
blasf77_dgemm( "No transpose", "No transpose",
&i__3, &i__4, &jb, &c_neg_one,
A(j+jb,j), &lda,
A(j,j+jb), &lda, &c_one,
A(j+jb,j+jb), &lda );
}
}
}
}
return *info;
} /* magma_dgetrf_nopiv */
/**
Purpose
-------
DPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
d_lA DOUBLE PRECISION array of pointers on the GPU, dimension (ngpu)
On entry, the symmetric matrix dA distributed over GPUs
(dl_A[d] points to the local matrix on the d-th GPU).
It is distributed in 1D block column or row cyclic (with the
block size of nb) if UPLO = MagmaUpper or MagmaLower, respectively.
If UPLO = MagmaUpper, the leading N-by-N upper triangular
part of dA contains the upper triangular part of the matrix dA,
and the strictly lower triangular part of dA is not referenced.
If UPLO = MagmaLower, the leading N-by-N lower triangular part
of dA contains the lower triangular part of the matrix dA, and
the strictly upper triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_dposv_comp
********************************************************************/
extern "C" magma_int_t
magma_dpotrf_mgpu_right(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n,
magmaDouble_ptr d_lA[], magma_int_t ldda,
magma_int_t *info )
{
#define dlA(id, i, j) (d_lA[(id)] + (j) * ldda + (i))
#define dlP(id, i, j) (d_lP[(id)] + (j) * ldda + (i))
#define panel(j) (panel + (j))
#define tmppanel(j) (tmppanel + (j))
#define tmpprevpanel(j) (tmpprevpanel + (j))
#define STREAM_ID(i) (nqueue > 1 ? 1+((i)/nb)%(nqueue-1) : 0)
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double d_one = 1.0;
double d_neg_one = -1.0;
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t j, nb, d, id, j_local, blkid, crosspoint, prevtrsmrows=0, nqueue = 5;
double *panel, *tmppanel0, *tmppanel1, *tmppanel, *tmpprevpanel;
double *d_lP[MagmaMaxGPUs], *dlpanel, *dlpanels[MagmaMaxGPUs];
magma_int_t rows, trsmrows, igpu, n_local[MagmaMaxGPUs], ldpanel;
magma_queue_t queues[MagmaMaxGPUs][10];
*info = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
nb = magma_get_dpotrf_nb(n);
ldpanel = ldda;
magma_setdevice(0);
if (MAGMA_SUCCESS != magma_dmalloc_pinned( &panel, 2 * nb * ldpanel )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
tmppanel0 = panel;
tmppanel1 = tmppanel0 + nb * ldpanel;
if ((nb <= 1) || (nb >= n)) {
// Use unblocked code.
magma_dgetmatrix( n, n, dlA(0, 0, 0), ldda, panel, ldpanel);
lapackf77_dpotrf( uplo_, &n, panel, &ldpanel, info);
magma_dsetmatrix( n, n, panel, ldpanel, dlA(0, 0, 0), ldda );
} else {
for( d = 0; d < ngpu; d++ ) {
// local-n and local-ld
n_local[d] = ((n / nb) / ngpu) * nb;
if (d < (n / nb) % ngpu)
n_local[d] += nb;
else if (d == (n / nb) % ngpu)
n_local[d] += n % nb;
magma_setdevice(d);
magma_device_sync();
if (MAGMA_SUCCESS != magma_dmalloc( &d_lP[d], nb * ldda )) {
for( j = 0; j < d; j++ ) {
magma_setdevice(j);
magma_free( d_lP[d] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
for( j=0; j < nqueue; j++ ) {
magma_queue_create( &queues[d][j] );
}
}
//#define ENABLE_TIMER
#if defined (ENABLE_TIMER)
real_Double_t therk[4], tmtc, tcchol, tctrsm, tctm, tmnp, tcnp;
real_Double_t ttot_herk[4] = {0,0,0,0}, ttot_mtc = 0, ttot_cchol = 0, ttot_ctrsm = 0, ttot_ctm = 0, ttot_mnp = 0, ttot_cnp = 0;
printf("\n\n %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s\n",
"j", "nb", "row", "mtc", "CPU_np", "panel", "ctrsm", "CH+TRSM", "CPU", "dsyrk[0]", "dsyrk[1]", "dsyrk[2]", "dsyrk[3]", "ctm P", "gpu_np");
printf(" ====================================================================================================\n");
#endif
// Use blocked code.
if (uplo == MagmaUpper) {
printf( " === not supported, yet ===\n" );
} else {
blkid = -1;
if (ngpu == 4)
crosspoint = n;
else if (ngpu == 3)
crosspoint = n;
else if (ngpu == 2)
crosspoint = 20160;
else
crosspoint = 0;
crosspoint = 0; //n; //n -- > gpu always does next panel, 0 --> cpu always does next panel
crosspoint = n;
#if defined (ENABLE_TIMER)
real_Double_t tget = magma_wtime(), tset = 0.0, ttot = 0.0;
#endif
if ( n > nb ) {
// send first panel to cpu
magma_setdevice(0);
tmppanel = tmppanel0;
magma_dgetmatrix_async(n, nb,
dlA(0, 0, 0), ldda,
tmppanel(0), ldpanel,
queues[0][0] );
}
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_device_sync();
}
tget = magma_wtime()-tget;
#endif
// Compute the Cholesky factorization A = L*L'
for (j = 0; (j + nb) < n; j += nb) {
#if defined (ENABLE_TIMER)
therk[0] = therk[1] = therk[2] = therk[3] = tmtc = tcchol = tctrsm = tctm = tmnp = tcnp = 0.0;
#endif
blkid += 1;
tmppanel = (blkid % 2 == 0) ? tmppanel0 : tmppanel1;
// Set the gpu number that holds the current panel
id = (j / nb) % ngpu;
magma_setdevice(id);
// Set the local index where the current panel is
j_local = j / (nb * ngpu) * nb;
rows = n - j;
// Wait for the panel on cpu
magma_queue_sync( queues[id][0] );
if (j > 0 && prevtrsmrows > crosspoint) {
#if defined (ENABLE_TIMER)
tcnp = magma_wtime();
#endif
tmpprevpanel = ((blkid - 1) % 2) == 0 ? tmppanel0 : tmppanel1;
blasf77_dgemm( MagmaNoTransStr, MagmaConjTransStr,
&rows, &nb, &nb,
&c_neg_one, tmpprevpanel(j), &ldpanel,
tmpprevpanel(j), &ldpanel,
&c_one, tmppanel(j), &ldpanel );
#if defined (ENABLE_TIMER)
tcnp = magma_wtime() - tcnp;
ttot_cnp += tcnp;
#endif
}
#if defined (ENABLE_TIMER)
tcchol = magma_wtime();
#endif
lapackf77_dpotrf(MagmaLowerStr, &nb, tmppanel(j), &ldpanel, info);
if (*info != 0) {
*info = *info + j;
break;
}
#if defined (ENABLE_TIMER)
tcchol = magma_wtime() - tcchol;
ttot_cchol += tcchol;
tctrsm = magma_wtime();
#endif
trsmrows = rows - nb;
if (trsmrows > 0) {
blasf77_dtrsm(MagmaRightStr, MagmaLowerStr, MagmaConjTransStr, MagmaNonUnitStr,
&trsmrows, &nb,
&c_one, tmppanel(j), &ldpanel,
tmppanel(j + nb), &ldpanel);
}
#if defined (ENABLE_TIMER)
tctrsm = magma_wtime() - tctrsm;
ttot_ctrsm += tctrsm;
tctm = magma_wtime();
#endif
d = (id + 1) % ngpu;
// send current panel to gpus
for (igpu = 0; igpu < ngpu; igpu++, d = (d + 1) % ngpu ) {
magma_int_t myrows = 0;
magma_int_t row_offset = 0;
if ( d == id ) {
dlpanel = dlA(d, j, j_local);
myrows = rows;
row_offset = 0;
} else {
dlpanel = dlP(d, 0, 0);
myrows = trsmrows;
row_offset = nb;
}
if (myrows > 0) {
magma_setdevice(d);
magma_dsetmatrix_async(myrows, nb,
tmppanel(j + row_offset), ldpanel,
dlpanel, ldda, queues[d][0] );
}
}
/* make sure panel is on GPUs */
d = (id + 1) % ngpu;
for (igpu = 0; igpu < ngpu; igpu++, d = (d + 1) % ngpu ) {
magma_setdevice(d);
magma_queue_sync( queues[d][0] );
}
#if defined (ENABLE_TIMER)
tctm = magma_wtime() - tctm;
ttot_ctm += tctm;
#endif
if ( (j + nb) < n) {
magma_int_t offset = 0;
magma_int_t row_offset = 0;
if (j + nb + nb < n) {
d = (id + 1) % ngpu;
magma_setdevice(d);
magma_int_t j_local2 = (j + nb) / (nb * ngpu) * nb;
if (trsmrows <= crosspoint) {
#if defined (ENABLE_TIMER)
tmnp = magma_wtime();
#endif
// do gemm on look ahead panel
if ( d == id ) {
dlpanel = dlA(d, j + nb, j_local);
} else {
dlpanel = dlP(d, 0, 0);
}
magmablasSetKernelStream( queues[d][STREAM_ID(j_local2)] );
#define DSYRK_ON_DIAG
#ifdef DSYRK_ON_DIAG
magma_dsyrk( MagmaLower, MagmaNoTrans,
nb, nb,
d_neg_one, dlpanel, ldda,
d_one, dlA(d, j + nb, j_local2), ldda);
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows-nb, nb, nb,
c_neg_one, dlpanel+nb, ldda,
dlpanel, ldda,
c_one, dlA(d, j + nb +nb, j_local2), ldda);
#else
magma_dgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows, nb, nb,
c_neg_one, dlpanel, ldda,
dlpanel, ldda,
c_one, dlA(d, j + nb, j_local2), ldda);
#endif
#if defined (ENABLE_TIMER)
magma_device_sync();
tmnp = magma_wtime() - tmnp;
ttot_mnp += tmnp;
#endif
}
// send next panel to cpu
magma_queue_sync( queues[d][STREAM_ID(j_local2)] ); // make sure lookahead is done
tmppanel = ((blkid+1) % 2 == 0) ? tmppanel0 : tmppanel1;
magma_dgetmatrix_async(rows-nb, nb,
dlA(d, j+nb, j_local2), ldda,
tmppanel(j+nb), ldpanel,
queues[d][0] );
tmppanel = (blkid % 2 == 0) ? tmppanel0 : tmppanel1;
offset = j + nb + nb;
row_offset = nb;
} else {
offset = j + nb;
row_offset = 0;
}
if (n - offset > 0) {
// syrk on multiple gpu
for (d = 0; d < ngpu; d++ ) {
if ( d == id ) {
dlpanels[d] = dlA(d, j + nb + row_offset, j_local);
} else {
dlpanels[d] = dlP(d, row_offset, 0);
}
}
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
therk[d] = magma_wtime();
}
#endif
//magmablasSetKernelStream( queues[d] );
//magma_dsyrk( MagmaLower, MagmaNoTrans, n - offset, nb,
// d_neg_one, dlpanel, ldda,
// d_one, &d_lA[d][offset + offset*ldda], ldda );
#ifdef DSYRK_ON_DIAG
magma_dsyrk_mgpu
#else
magma_dsyrk_mgpu2
#endif
(ngpu, MagmaLower, MagmaNoTrans,
nb, n - offset, nb,
d_neg_one, dlpanels, ldda, 0,
d_one, d_lA, ldda, offset,
nqueue, queues );
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_device_sync();
therk[d] = magma_wtime() - therk[d];
ttot_herk[d] += therk[d];
}
#endif
}
prevtrsmrows = trsmrows;
#if defined (ENABLE_TIMER)
ttot += (tcnp+tcchol+tctrsm+therk[0]+therk[1]+therk[2]+tctm+tmnp);
printf("%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(%d) %10.3lf\n",
j, nb, rows, tmtc,
tcnp, // gemm
tcchol, // potrf
tctrsm, // trsm
(tcchol + tctrsm),
(tmtc+tcnp+tcchol+tctrsm),
therk[0], therk[1], therk[2], therk[3], // syrk
tctm, // copy panel to GPU
tmnp, // lookahead on GPU
(id + 1) % ngpu,
(tcnp+tcchol+tctrsm+therk[0]+therk[1]+therk[2]+tctm+tmnp));
fflush(0);
#endif
}
}
for( d = 0; d < ngpu; d++ ) {
magma_setdevice(d);
for( id=0; id < nqueue; id++ ) {
magma_queue_sync( queues[d][id] );
}
}
#if defined (ENABLE_TIMER)
printf("\n%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(-) %10.3lf\n",
n, n, 0, ttot_mtc,
ttot_cnp, // gemm
ttot_cchol, // potrf
ttot_ctrsm, // trsm
(ttot_cchol + ttot_ctrsm),
(ttot_mtc+ttot_cnp+ttot_cchol+ttot_ctrsm),
ttot_herk[0], ttot_herk[1], ttot_herk[2], ttot_herk[3], // syrk
ttot_ctm, // copy panel to GPU
ttot_mnp, // lookahead on GPU
(ttot_cnp+ttot_cchol+ttot_ctrsm+ttot_herk[0]+ttot_herk[1]+ttot_herk[2]+ttot_ctm+ttot_mnp));
printf("%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(-) %10.3lf (ratio)\n",
n, n, 0, ttot_mtc/ttot,
ttot_cnp/ttot, // gemm
ttot_cchol/ttot, // potrf
ttot_ctrsm/ttot, // trsm
(ttot_cchol + ttot_ctrsm)/ttot,
(ttot_mtc+ttot_cnp+ttot_cchol+ttot_ctrsm)/ttot,
ttot_herk[0]/ttot, ttot_herk[1]/ttot, ttot_herk[2]/ttot, ttot_herk[3]/ttot, // syrk
ttot_ctm/ttot, // copy panel to GPU
ttot_mnp/ttot, // lookahead on GPU
(ttot_cnp+ttot_cchol+ttot_ctrsm+ttot_herk[0]+ttot_herk[1]+ttot_herk[2]+ttot_ctm+ttot_mnp)/ttot);
#endif
// cholesky for the last block
if (j < n && *info == 0) {
rows = n - j;
id = (j / nb) % ngpu;
// Set the local index where the current panel is
j_local = j / (nb * ngpu) * nb;
magma_setdevice(id);
#if defined (ENABLE_TIMER)
tset = magma_wtime();
#endif
magma_dgetmatrix(rows, rows, dlA(id, j, j_local), ldda, panel(j), ldpanel);
lapackf77_dpotrf(MagmaLowerStr, &rows, panel(j), &ldpanel, info);
magma_dsetmatrix(rows, rows, panel(j), ldpanel, dlA(id, j, j_local), ldda);
#if defined (ENABLE_TIMER)
tset = magma_wtime() - tset;
#endif
}
#if defined (ENABLE_TIMER)
printf( " matrix_get,set: %10.3lf %10.3lf -> %10.3lf\n",tget,tset,ttot+tget+tset );
#endif
} // end of else not upper
// clean up
for( d = 0; d < ngpu; d++ ) {
magma_setdevice(d);
for( j=0; j < nqueue; j++ ) {
magma_queue_destroy( queues[d][j] );
}
magma_free( d_lP[d] );
}
} // end of not lapack
// free workspace
magma_free_pinned( panel );
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_dpotrf_mgpu_right */