/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgetrf
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double *h_A, *h_R, *work;
magmaDouble_ptr d_A, dwork;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t N, n2, lda, ldda, info, lwork, ldwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double tmp;
double error, rwork[1];
magma_int_t *ipiv;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// need looser bound (3000*eps instead of 30*eps) for tests
// TODO: should compute ||I - A*A^{-1}|| / (n*||A||*||A^{-1}||)
opts.tolerance = max( 3000., opts.tolerance );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf(" N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / (N*||A||_F)\n");
printf("=================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = ((N+31)/32)*32;
ldwork = N * magma_get_dgetri_nb( N );
gflops = FLOPS_DGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0)
printf("lapackf77_dgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
lwork = int( MAGMA_D_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, double, lwork );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( dwork, double, ldwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
error = lapackf77_dlange( "f", &N, &N, h_A, &lda, rwork ); // norm(A)
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_dsetmatrix( N, N, h_A, lda, d_A, 0, ldda, opts.queue );
magma_dgetrf_gpu( N, N, d_A, 0, ldda, ipiv, opts.queue, &info );
magma_dgetmatrix( N, N, d_A, 0, ldda, h_A, lda, opts.queue );
if ( info != 0 )
printf("magma_dgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// check for exact singularity
//h_A[ 10 + 10*lda ] = MAGMA_D_MAKE( 0.0, 0.0 );
//magma_dsetmatrix( N, N, h_A, lda, d_A, 0, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dgetri_gpu( N, d_A, 0, ldda, ipiv, dwork, 0, ldwork, opts.queues2, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix( N, N, d_A, 0, ldda, h_R, lda, opts.queue );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgetri( &N, h_A, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_daxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_dlange( "f", &N, &N, h_R, &lda, rwork ) / (N*error);
printf( "%5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
printf( "%5d --- ( --- ) %7.2f (%7.2f) ---\n",
(int) N, gpu_perf, gpu_time );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgetri
*/
int main( int argc, char** argv )
{
TESTING_INIT();
// constants
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double *h_A, *h_Ainv, *h_R, *work;
magmaDouble_ptr d_A, dwork;
magma_int_t N, n2, lda, ldda, info, lwork, ldwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
double tmp;
double error, rwork[1];
magma_int_t *ipiv;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% N CPU Gflop/s (sec) GPU Gflop/s (sec) ||I - A*A^{-1}||_1 / (N*cond(A))\n");
printf("%%===============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
ldwork = N * magma_get_dgetri_nb( N );
gflops = FLOPS_DGETRI( N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0) {
printf("lapackf77_dgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = magma_int_t( MAGMA_D_REAL( tmp ));
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
TESTING_MALLOC_CPU( work, double, lwork );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_Ainv, double, n2 );
TESTING_MALLOC_CPU( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( dwork, double, ldwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
/* Factor the matrix. Both MAGMA and LAPACK will use this factor. */
magma_dsetmatrix( N, N, h_A, lda, d_A, ldda, opts.queue );
magma_dgetrf_gpu( N, N, d_A, ldda, ipiv, &info );
magma_dgetmatrix( N, N, d_A, ldda, h_Ainv, lda, opts.queue );
if (info != 0) {
printf("magma_dgetrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
// check for exact singularity
//h_Ainv[ 10 + 10*lda ] = MAGMA_D_MAKE( 0.0, 0.0 );
//magma_dsetmatrix( N, N, h_Ainv, lda, d_A, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_dgetri_gpu( N, d_A, ldda, ipiv, dwork, ldwork, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_dgetri_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgetri( &N, h_Ainv, &lda, ipiv, work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_dgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
printf( "%5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
}
else {
printf( "%5d --- ( --- ) %7.2f (%7.2f)",
(int) N, gpu_perf, gpu_time );
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.check ) {
magma_dgetmatrix( N, N, d_A, ldda, h_Ainv, lda, opts.queue );
// compute 1-norm condition number estimate, following LAPACK's zget03
double normA, normAinv, rcond;
normA = lapackf77_dlange( "1", &N, &N, h_A, &lda, rwork );
normAinv = lapackf77_dlange( "1", &N, &N, h_Ainv, &lda, rwork );
if ( normA <= 0 || normAinv <= 0 ) {
rcond = 0;
error = 1 / (tol/opts.tolerance); // == 1/eps
}
else {
rcond = (1 / normA) / normAinv;
// R = I
// R -= A*A^{-1}
// err = ||I - A*A^{-1}|| / ( N ||A||*||A^{-1}|| ) = ||R|| * rcond / N, using 1-norm
lapackf77_dlaset( "full", &N, &N, &c_zero, &c_one, h_R, &lda );
blasf77_dgemm( "no", "no", &N, &N, &N,
&c_neg_one, h_A, &lda,
h_Ainv, &lda,
&c_one, h_R, &lda );
error = lapackf77_dlange( "1", &N, &N, h_R, &lda, rwork );
error = error * rcond / N;
}
bool okay = (error < tol);
status += ! okay;
printf( " %8.2e %s\n",
error, (okay ? "ok" : "failed"));
}
else {
printf( "\n" );
}
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_Ainv );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
magma_int_t magmaf_get_dgetri_nb( magma_int_t *m )
{
return magma_get_dgetri_nb( *m );
}
extern "C" magma_int_t
magma_dgetri_gpu( magma_int_t n, double *dA, magma_int_t lda,
magma_int_t *ipiv, double *dwork, magma_int_t lwork,
magma_int_t *info )
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
DGETRI computes the inverse of a matrix using the LU factorization
computed by DGETRF. This method inverts U and then computes inv(A) by
solving the system inv(A)*L = inv(U) for inv(A).
Note that it is generally both faster and more accurate to use DGESV,
or DGETRF and DGETRS, to solve the system AX = B, rather than inverting
the matrix and multiplying to form X = inv(A)*B. Only in special
instances should an explicit inverse be computed with this routine.
Arguments
=========
N (input) INTEGER
The order of the matrix A. N >= 0.
dA (input/output) DOUBLE_PRECISION array on the GPU, dimension (LDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by DGETRF_GPU.
On exit, if INFO = 0, the inverse of the original matrix A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from DGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
DWORK (workspace/output) COMPLEX*16 array on the GPU, dimension (MAX(1,LWORK))
LWORK (input) INTEGER
The dimension of the array DWORK. LWORK >= N*NB, where NB is
the optimal blocksize returned by magma_get_dgetri_nb(n).
Unlike LAPACK, this version does not currently support a
workspace query, because the workspace is on the GPU.
INFO (output) 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 matrix is
singular and its cannot be computed.
===================================================================== */
/* Local variables */
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *dL = dwork;
magma_int_t ldl = n;
magma_int_t nb = magma_get_dgetri_nb(n);
magma_int_t j, jmax, jb, jp;
*info = 0;
if (n < 0)
*info = -1;
else if (lda < max(1,n))
*info = -3;
else if ( lwork < n*nb )
*info = -6;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( n == 0 )
return *info;
/* Invert the triangular factor U */
magma_dtrtri_gpu( MagmaUpper, MagmaNonUnit, n, dA, lda, info );
if ( *info != 0 )
return *info;
jmax = ((n-1) / nb)*nb;
for( j = jmax; j >= 0; j -= nb ) {
jb = min( nb, n-j );
// copy current block column of L to work space,
// then replace with zeros in A.
magmablas_dlacpy( MagmaUpperLower, n-j, jb,
&dA[j + j*lda], lda,
&dL[j ], ldl );
magmablas_dlaset( MagmaLower, n-j, jb, &dA[j + j*lda], lda );
// compute current block column of Ainv
// Ainv(:, j:j+jb-1)
// = ( U(:, j:j+jb-1) - Ainv(:, j+jb:n) L(j+jb:n, j:j+jb-1) )
// * L(j:j+jb-1, j:j+jb-1)^{-1}
// where L(:, j:j+jb-1) is stored in dL.
if ( j+jb < n ) {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, n, jb, n-j-jb,
c_neg_one, &dA[(j+jb)*lda], lda,
&dL[ j+jb ], ldl,
c_one, &dA[ j*lda], lda );
}
magma_dtrsm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaUnit,
n, jb, c_one,
&dL[j ], ldl,
&dA[j*lda], lda );
}
// Apply column interchanges
for( j = n-2; j >= 0; --j ) {
jp = ipiv[j] - 1;
if ( jp != j ) {
magmablas_dswap( n, &dA[ j*lda ], 1, &dA[ jp*lda ], 1 );
}
}
return *info;
}
SEXP magma_dgeMatrix_solve(SEXP a)
{
#ifdef HIPLAR_WITH_MAGMA
/* compute the 1-norm of the matrix, which is needed
later for the computation of the reciprocal condition number. */
double aNorm = magma_get_norm(a, "1");
/* the LU decomposition : */
/* Given that we may be performing this operation
* on the GPU we may put in an optimisation here
* where if we call the LU solver we, we do not require
* the decomposition to be transferred back to CPU. This is TODO
*/
SEXP val = PROTECT(NEW_OBJECT(MAKE_CLASS("dgeMatrix"))),
lu = magma_dgeMatrix_LU_(a, TRUE);
int *dims = INTEGER(GET_SLOT(lu, Matrix_DimSym)),
*pivot = INTEGER(GET_SLOT(lu, Matrix_permSym));
/* prepare variables for the dgetri calls */
double *x, tmp;
int info, lwork = -1;
if (dims[0] != dims[1]) error(_("Solve requires a square matrix"));
slot_dup(val, lu, Matrix_xSym);
x = REAL(GET_SLOT(val, Matrix_xSym));
slot_dup(val, lu, Matrix_DimSym);
int N2 = dims[0] * dims[0];
if(dims[0]) /* the dimension is not zero */
{
/* is the matrix is *computationally* singular ? */
double rcond;
F77_CALL(dgecon)("1", dims, x, dims, &aNorm, &rcond,
(double *) R_alloc(4*dims[0], sizeof(double)),
(int *) R_alloc(dims[0], sizeof(int)), &info);
if (info)
error(_("error [%d] from Lapack 'dgecon()'"), info);
if(rcond < DOUBLE_EPS)
error(_("Lapack dgecon(): system computationally singular, reciprocal condition number = %g"),
rcond);
/* only now try the inversion and check if the matrix is *exactly* singular: */
// This is also a work space query. This is not an option in magma
F77_CALL(dgetri)(dims, x, dims, pivot, &tmp, &lwork, &info);
lwork = (int) tmp;
if( GPUFlag == 0){
F77_CALL(dgetri)(dims, x, dims, pivot,
(double *) R_alloc((size_t) lwork, sizeof(double)),
&lwork, &info);
#ifdef HIPLAR_DBG
R_ShowMessage("DBG: Solve using LU using dgetri;");
#endif
}
else if(GPUFlag == 1) {
double *d_x, *dwork;
cublasStatus retStatus;
#ifdef HIPLAR_DBG
R_ShowMessage("Solve using LU using magma_dgetri;");
#endif
cublasAlloc(N2, sizeof(double), (void**)&d_x);
//cublasAlloc(N2 , sizeof(double), (void**)&dtmp);
/* Error Checking */
retStatus = cublasGetError ();
if (retStatus != CUBLAS_STATUS_SUCCESS)
error(_("CUBLAS: Error in Memory Allocation on Device"));
/********************************************/
cublasSetVector( N2, sizeof(double), x, 1, d_x, 1);
/* Error Checking */
retStatus = cublasGetError ();
if (retStatus != CUBLAS_STATUS_SUCCESS)
error(_("CUBLAS: Error in Data Transfer to Device"));
/********************************************/
lwork = dims[0] * magma_get_dgetri_nb( dims[0] );
cublasAlloc(lwork, sizeof(double), (void**)&dwork);
/* Error Checking */
retStatus = cublasGetError ();
if (retStatus != CUBLAS_STATUS_SUCCESS)
error(_("CUBLAS: Error in Memory Allocation on Device"));
/********************************************/
magma_dgetri_gpu(dims[0], d_x, dims[0], pivot, dwork , lwork, &info);
cublasGetVector(N2, sizeof(double), d_x, 1, x, 1);
/* Error Checking */
retStatus = cublasGetError ();
if (retStatus != CUBLAS_STATUS_SUCCESS)
error(_("CUBLAS: Error in Data From to Device"));
/********************************************/
cublasFree(dwork);
cublasFree(d_x);
/* Error Checking */
retStatus = cublasGetError ();
if (retStatus != CUBLAS_STATUS_SUCCESS)
error(_("CUBLAS: Error freeing memory"));
/********************************************/
}
else
error(_("GPUFlag not set correctly"));
if (info)
error(_("Lapack routine dgetri: system is exactly singular"));
}
UNPROTECT(1);
return val;
#endif
return R_NilValue;
}
/**
Purpose
-------
DGETRI computes the inverse of a matrix using the LU factorization
computed by DGETRF. This method inverts U and then computes inv(A) by
solving the system inv(A)*L = inv(U) for inv(A).
Note that it is generally both faster and more accurate to use DGESV,
or DGETRF and DGETRS, to solve the system AX = B, rather than inverting
the matrix and multiplying to form X = inv(A)*B. Only in special
instances should an explicit inverse be computed with this routine.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA DOUBLE_PRECISION array on the GPU, dimension (LDDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by DGETRF_GPU.
On exit, if INFO = 0, the inverse of the original matrix A.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[in]
ipiv INTEGER array, dimension (N)
The pivot indices from DGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[out]
dwork (workspace) DOUBLE_PRECISION array on the GPU, dimension (MAX(1,LWORK))
@param[in]
lwork INTEGER
The dimension of the array DWORK. LWORK >= N*NB, where NB is
the optimal blocksize returned by magma_get_dgetri_nb(n).
\n
Unlike LAPACK, this version does not currently support a
workspace query, because the workspace is on the GPU.
@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 matrix is
singular and its cannot be computed.
@ingroup magma_dgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_dgetri_gpu( magma_int_t n, double *dA, magma_int_t ldda,
magma_int_t *ipiv, double *dwork, magma_int_t lwork,
magma_int_t *info )
{
#define dA(i, j) (dA + (i) + (j)*ldda)
#define dL(i, j) (dL + (i) + (j)*lddl)
/* Local variables */
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *dL = dwork;
magma_int_t lddl = n;
magma_int_t nb = magma_get_dgetri_nb(n);
magma_int_t j, jmax, jb, jp;
*info = 0;
if (n < 0)
*info = -1;
else if (ldda < max(1,n))
*info = -3;
else if ( lwork < n*nb )
*info = -6;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( n == 0 )
return *info;
/* Invert the triangular factor U */
magma_dtrtri_gpu( MagmaUpper, MagmaNonUnit, n, dA, ldda, info );
if ( *info != 0 )
return *info;
jmax = ((n-1) / nb)*nb;
for( j = jmax; j >= 0; j -= nb ) {
jb = min( nb, n-j );
// copy current block column of A to work space dL
// (only needs lower trapezoid, but we also copy upper triangle),
// then zero the strictly lower trapezoid block column of A.
magmablas_dlacpy( MagmaFull, n-j, jb,
dA(j,j), ldda,
dL(j,0), lddl );
magmablas_dlaset( MagmaLower, n-j-1, jb, c_zero, c_zero, dA(j+1,j), ldda );
// compute current block column of Ainv
// Ainv(:, j:j+jb-1)
// = ( U(:, j:j+jb-1) - Ainv(:, j+jb:n) L(j+jb:n, j:j+jb-1) )
// * L(j:j+jb-1, j:j+jb-1)^{-1}
// where L(:, j:j+jb-1) is stored in dL.
if ( j+jb < n ) {
magma_dgemm( MagmaNoTrans, MagmaNoTrans, n, jb, n-j-jb,
c_neg_one, dA(0,j+jb), ldda,
dL(j+jb,0), lddl,
c_one, dA(0,j), ldda );
}
// TODO use magmablas work interface
magma_dtrsm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaUnit,
n, jb, c_one,
dL(j,0), lddl,
dA(0,j), ldda );
}
// Apply column interchanges
for( j = n-2; j >= 0; --j ) {
jp = ipiv[j] - 1;
if ( jp != j ) {
magmablas_dswap( n, dA(0,j), 1, dA(0,jp), 1 );
}
}
return *info;
}
