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

Go to the source code of this file.

Macros

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

Functions

magma_int_t magma_ztrtri_gpu (char uplo, char diag, magma_int_t n, cuDoubleComplex *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 ztrtri_gpu.cpp.

#define PRECISION_z

Definition at line 14 of file ztrtri_gpu.cpp.

Function Documentation

magma_int_t magma_ztrtri_gpu ( char  uplo,
char  diag,
magma_int_t  n,
cuDoubleComplex *  dA,
magma_int_t  ldda,
magma_int_t *  info 
)

Definition at line 31 of file ztrtri_gpu.cpp.

References __func__, dA, diag, lapackf77_lsame, lapackf77_ztrtri, MAGMA_ERR_HOST_ALLOC, magma_free_host(), magma_get_zpotrf_nb(), magma_queue_create(), magma_queue_destroy(), magma_queue_sync(), MAGMA_SUCCESS, magma_xerbla(), MAGMA_Z_NEG_ONE, MAGMA_Z_ONE, magma_zgetmatrix(), magma_zgetmatrix_async(), magma_zmalloc_host(), magma_zsetmatrix(), magma_zsetmatrix_async(), magma_ztrmm(), magma_ztrsm(), 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;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
cuDoubleComplex *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_zpotrf_nb(n);
if (MAGMA_SUCCESS != magma_zmalloc_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_zgetmatrix( n, n, dA, ldda, work, n );
lapackf77_ztrtri(uplo_, diag_, &n, work, &n, info);
magma_zsetmatrix( 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_ztrmm(MagmaLeft, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_one, dA(0,0), ldda, dA(0, j),ldda);
magma_ztrsm(MagmaRight, MagmaUpper,
MagmaNoTrans, MagmaNonUnit, j, jb,
c_neg_one, dA(j,j), ldda, dA(0, j),ldda);
//cublasGetMatrix(jb ,jb, sizeof(cuDoubleComplex),
// dA(j, j), ldda, work, jb);
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ztrtri(MagmaUpperStr, diag_, &jb, work, &jb, info);
//cublasSetMatrix(jb, jb, sizeof(cuDoubleComplex),
// work, jb, dA(j, j), ldda);
magma_zsetmatrix_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_ztrmm(MagmaLeft, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_one, dA(j+jb,j+jb), ldda, dA(j+jb, j), ldda);
magma_ztrsm(MagmaRight, MagmaLower,
MagmaNoTrans, MagmaNonUnit, (n-j-jb), jb,
c_neg_one, dA(j,j), ldda, dA(j+jb, j), ldda);
}
//cublasGetMatrix(jb, jb, sizeof(cuDoubleComplex),
// dA(j, j), ldda, work, jb);
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, stream[1] );
magma_queue_sync( stream[1] );
/* Compute inverse of current diagonal block */
lapackf77_ztrtri(MagmaLowerStr, diag_, &jb, work, &jb, info);
//cublasSetMatrix(jb, jb, sizeof(cuDoubleComplex),
// work, jb, dA(j, j), ldda);
magma_zsetmatrix_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: