MAGMA  1.2.0
MatrixAlgebraonGPUandMulticoreArchitectures
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups
Macros | Functions
dtrtri_gpu.cpp File Reference
#include "common_magma.h"
Include dependency graph for dtrtri_gpu.cpp:

Go to the source code of this file.

Macros

#define PRECISION_d
#define magma_dgemm   magmablas_dgemm
#define magma_dtrsm   magmablas_dtrsm
#define dA(i, j)   (dA+(j)*ldda + (i))

Functions

magma_int_t magma_dtrtri_gpu (char uplo, char diag, magma_int_t n, double *dA, magma_int_t ldda, magma_int_t *info)

Macro Definition Documentation

#define dA (   i,
  j 
)    (dA+(j)*ldda + (i))

Definition at line 28 of file dtrtri_gpu.cpp.

#define magma_dgemm   magmablas_dgemm

Definition at line 16 of file dtrtri_gpu.cpp.

#define magma_dtrsm   magmablas_dtrsm

Definition at line 17 of file dtrtri_gpu.cpp.

#define PRECISION_d

Definition at line 14 of file dtrtri_gpu.cpp.

Function Documentation

magma_int_t magma_dtrtri_gpu ( char  uplo,
char  diag,
magma_int_t  n,
double *  dA,
magma_int_t  ldda,
magma_int_t *  info 
)

Definition at line 31 of file dtrtri_gpu.cpp.

References __func__, dA, diag, lapackf77_dtrtri(), lapackf77_lsame, MAGMA_D_NEG_ONE, MAGMA_D_ONE, magma_dgetmatrix(), magma_dgetmatrix_async(), magma_dmalloc_host(), magma_dsetmatrix(), magma_dsetmatrix_async(), magma_dtrmm(), magma_dtrsm, MAGMA_ERR_HOST_ALLOC, magma_free_host(), magma_get_dpotrf_nb(), magma_queue_create(), magma_queue_destroy(), magma_queue_sync(), MAGMA_SUCCESS, magma_xerbla(), MagmaLeft, MagmaLower, MagmaLowerStr, MagmaNonUnit, MagmaNoTrans, MagmaRight, MagmaUpper, MagmaUpperStr, max, min, uplo, and codegen::work.

{
/* -- MAGMA (version 1.2.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
May 2012
Purpose
=======
DTRTRI computes the inverse of a real upper or lower triangular
matrix dA.
This is the Level 3 BLAS version of the algorithm.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': A is upper triangular;
= 'L': A is lower triangular.
DIAG (input) CHARACTER*1
= 'N': A is non-unit triangular;
= 'U': A is unit triangular.
N (input) INTEGER
The order of the matrix A. N >= 0.
dA (input/output) DOUBLE PRECISION array ON THE GPU, dimension (LDDA,N)
On entry, the triangular matrix A. If UPLO = 'U', the
leading N-by-N upper triangular part of the array dA contains
the upper triangular matrix, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of the array dA contains
the lower triangular matrix, and the strictly upper
triangular part of A is not referenced. If DIAG = 'U', the
diagonal elements of A are also not referenced and are
assumed to be 1.
On exit, the (triangular) inverse of the original matrix, in
the same storage format.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, dA(i,i) is exactly zero. The triangular
matrix is singular and its inverse can not be computed.
===================================================================== */
/* Local variables */
char uplo_[2] = {uplo, 0};
char diag_[2] = {diag, 0};
magma_int_t nb, nn, j, jb;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double *work;
long int upper = lapackf77_lsame(uplo_, "U");
long int nounit = lapackf77_lsame(diag_, "N");
*info = 0;
if ((! upper) && (! lapackf77_lsame(uplo_, "L")))
*info = -1;
else if ((! nounit) && (! lapackf77_lsame(diag_, "U")))
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Check for singularity if non-unit */
if (nounit)
{
for (*info=0; *info < n; *info=*info+1)
{
if(dA(*info,*info)==0)
return *info;
}
*info=0;
}
nb = magma_get_dpotrf_nb(n);
if (MAGMA_SUCCESS != magma_dmalloc_host( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
static cudaStream_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n)
{
magma_dgetmatrix( n, n, dA, ldda, work, n );
lapackf77_dtrtri(uplo_, diag_, &n, work, &n, info);
magma_dsetmatrix( n, n, work, n, dA, ldda );
}
else
{
if (upper)
{
/* Compute inverse of upper triangular matrix */
for (j=0; j<n; j =j+ nb)
{
jb = min(nb, (n-j));
/* Compute rows 1:j-1 of current block column */
magma_dtrmm(MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_dtrsm(MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
//cublasGetMatrix(jb ,jb, sizeof(double),
// dA(j, j), ldda, work, jb);
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_dtrtri(MagmaUpperStr, diag_, &jb, work, &jb, info);
//cublasSetMatrix(jb, jb, sizeof(double),
// work, jb, dA(j, j), ldda);
magma_dsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[0] );
}
}
else
{
/* Compute inverse of lower triangular matrix */
nn=((n-1)/nb)*nb+1;
for(j=nn-1; j>=0; j=j-nb)
{
jb=min(nb,(n-j));
if((j+jb) < n)
{
/* Compute rows j+jb:n of current block column */
magma_dtrmm(MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda);
magma_dtrsm(MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda);
}
//cublasGetMatrix(jb, jb, sizeof(double),
// dA(j, j), ldda, work, jb);
magma_dgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_dtrtri(MagmaLowerStr, diag_, &jb, work, &jb, info);
//cublasSetMatrix(jb, jb, sizeof(double),
// work, jb, dA(j, j), ldda);
magma_dsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, stream[0] );
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_host( work );
return *info;
}

Here is the call graph for this function:

Here is the caller graph for this function: