/**
Purpose
-------
ZPOTRS solves a system of linear equations A*X = B with a Hermitian
positive definite matrix A using the Cholesky factorization
A = U**H*U or A = L*L**H computed by ZPOTRF.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
The triangular factor U or L from the Cholesky factorization
A = U**H*U or A = L*L**H, as computed by ZPOTRF.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[in,out]
dB COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zposv_comp
********************************************************************/
extern "C" magma_int_t
magma_zpotrs_gpu(
magma_uplo_t uplo, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
magmaDoubleComplex c_one = MAGMA_Z_ONE;
*info = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower )
*info = -1;
if ( n < 0 )
*info = -2;
if ( nrhs < 0)
*info = -3;
if ( ldda < max(1, n) )
*info = -5;
if ( lddb < max(1, n) )
*info = -7;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( (n == 0) || (nrhs == 0) ) {
return *info;
}
if ( uplo == MagmaUpper ) {
if ( nrhs == 1) {
magma_ztrsv(MagmaUpper, MagmaConjTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
}
}
else {
if ( nrhs == 1) {
magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaLower, MagmaConjTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
}
}
return *info;
}
/**
Purpose
-------
Solves a system of linear equations
A * X = B, A**T * X = B, or A**H * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in]
dA COMPLEX_16 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in]
ipiv INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[in,out]
dB COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrs_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, magma_int_t ldda, magma_int_t *ipiv,
magmaDoubleComplex_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex *work = NULL;
int notran = (trans == MagmaNoTrans);
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(trans != MagmaTrans) &&
(trans != MagmaConjTrans) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (n == 0 || nrhs == 0) {
return *info;
}
magma_zmalloc_cpu( &work, n * nrhs );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
i1 = 1;
i2 = n;
if (notran) {
inc = 1;
/* Solve A * X = B. */
magma_zgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, n, dB, lddb );
if ( nrhs == 1) {
magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
} else {
inc = -1;
/* Solve A**T * X = B or A**H * X = B. */
if ( nrhs == 1) {
magma_ztrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaLower, trans, MagmaUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ztrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
magma_zgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, n, dB, lddb );
}
magma_free_cpu(work);
return *info;
}
/**
Purpose
-------
ZGETRS solves a system of linear equations
A * X = B,
A**T * X = B, or
A**H * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_NOPIV_GPU.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[in,out]
dB COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrs_nopiv_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
// Constants
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
// Local variables
bool notran = (trans == MagmaNoTrans);
*info = 0;
if ( (! notran) &&
(trans != MagmaTrans) &&
(trans != MagmaConjTrans) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (n == 0 || nrhs == 0) {
return *info;
}
magma_queue_t queue = NULL;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
if (notran) {
/* Solve A * X = B. */
if ( nrhs == 1) {
magma_ztrsv( MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, ldda, dB, 1, queue );
magma_ztrsv( MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1, queue );
} else {
magma_ztrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb, queue );
magma_ztrsm( MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb, queue );
}
} else {
/* Solve A**T * X = B or A**H * X = B. */
if ( nrhs == 1) {
magma_ztrsv( MagmaUpper, trans, MagmaNonUnit, n, dA, ldda, dB, 1, queue );
magma_ztrsv( MagmaLower, trans, MagmaUnit, n, dA, ldda, dB, 1, queue );
} else {
magma_ztrsm( MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb, queue );
magma_ztrsm( MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb, queue );
}
}
magma_queue_destroy( queue );
return *info;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaDoubleComplex *A, *B, *C, *C2, *LU;
magmaDoubleComplex *dA, *dB, *dC1, *dC2;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.5, 0.1 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_int_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
printf("=========================================================================\n");
printf( "m=%d, n=%d, k=%d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_zmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_zmalloc( &dA, size ); assert( err == 0 );
err = magma_zmalloc( &dB, size ); assert( err == 0 );
err = magma_zmalloc( &dC1, size ); assert( err == 0 );
err = magma_zmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_zlarnv( &ione, ISEED, &size, A );
lapackf77_zlarnv( &ione, ISEED, &size, B );
lapackf77_zlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test ZSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetmatrix( m, n, A, ld, dB, ld );
magma_zswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_zswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_zgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "zswap diff %.2g\n", error );
}
else {
printf( "zswap skipped for n < 3\n" );
}
// ----- test IZAMAX
// get argmax of column of A
magma_zsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_izamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIzamax( handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "izamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test ZGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetvector( maxn, B, 1, dB, 1 );
magma_zsetvector( maxn, C, 1, dC1, 1 );
magma_zsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemv( handle, cublas_trans_const(trans[ia]),
m, n, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == MagmaNoTrans ? m : n);
cublasZaxpy( handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMV( m, n ) / 1e9;
printf( "zgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZHEMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetvector( m, B, 1, dB, 1 );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zhemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemv( handle, cublas_uplo_const(uplo[iu]),
m, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMV( m ) / 1e9;
printf( "zhemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_zlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_zgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_zsetmatrix( m, m, LU, ld, dA, ld );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ztrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsv( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ztrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test ZGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_zsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_zsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemm( handle, cublas_trans_const(trans[ia]), cublas_trans_const(trans[ib]),
m, n, k, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMM( m, n, k ) / 1e9;
printf( "zgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHEMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetmatrix( m, n, B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zhemm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
m, n, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMM( side[is], m, n ) / 1e9;
printf( "zhemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHERK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_zsetmatrix( n, k, A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZherk( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &dalpha, dA, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHERK( k, n ) / 1e9;
printf( "zherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHER2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == MagmaNoTrans);
magma_zsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zher2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZher2k( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &alpha, dA, ld, dB, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHER2K( k, n ) / 1e9;
printf( "zher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasZtrmm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRMM( side[is], m, n ) / 1e9;
printf( "ztrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test ZTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( side[is], m, n ) / 1e9;
printf( "ztrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
fflush( stdout );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
int status = (total_error != 0.);
return status;
}
extern "C" magma_int_t
magma_zpidr_merge(
magma_z_matrix A, magma_z_matrix b, magma_z_matrix *x,
magma_z_solver_par *solver_par,
magma_z_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PIDRMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->init_res = 0.0;
solver_par->final_res = 0.0;
solver_par->iter_res = 0.0;
solver_par->runtime = 0.0;
// constants
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_n_one = MAGMA_Z_NEG_ONE;
// internal user parameters
const magma_int_t smoothing = 1; // 0 = disable, 1 = enable
const double angle = 0.7; // [0-1]
// local variables
magma_int_t iseed[4] = {0, 0, 0, 1};
magma_int_t dof;
magma_int_t s;
magma_int_t distr;
magma_int_t k, i, sk;
magma_int_t innerflag;
magma_int_t ldd;
double residual;
double nrm;
double nrmb;
double nrmr;
double nrmt;
double rho;
magmaDoubleComplex om;
magmaDoubleComplex gamma;
magmaDoubleComplex fk;
// matrices and vectors
magma_z_matrix dxs = {Magma_CSR};
magma_z_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_z_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_z_matrix dG = {Magma_CSR}, dGcol = {Magma_CSR};
magma_z_matrix dU = {Magma_CSR};
magma_z_matrix dM = {Magma_CSR}, hMdiag = {Magma_CSR};
magma_z_matrix df = {Magma_CSR};
magma_z_matrix dt = {Magma_CSR}, dtt = {Magma_CSR};
magma_z_matrix dc = {Magma_CSR};
magma_z_matrix dv = {Magma_CSR};
magma_z_matrix dlu = {Magma_CSR};
magma_z_matrix dskp = {Magma_CSR}, hskp = {Magma_CSR};
magma_z_matrix dalpha = {Magma_CSR}, halpha = {Magma_CSR};
magma_z_matrix dbeta = {Magma_CSR}, hbeta = {Magma_CSR};
magmaDoubleComplex *d1 = NULL, *d2 = NULL;
// chronometry
real_Double_t tempo1, tempo2;
// initial s space
// TODO: add option for 's' (shadow space number)
// Hack: uses '--restart' option as the shadow space number.
// This is not a good idea because the default value of restart option is used to detect
// if the user provided a custom restart. This means that if the default restart value
// is changed then the code will think it was the user (unless the default value is
// also updated in the 'if' statement below.
s = 1;
if ( solver_par->restart != 50 ) {
if ( solver_par->restart > A.num_cols ) {
s = A.num_cols;
} else {
s = solver_par->restart;
}
}
solver_par->restart = s;
// set max iterations
solver_par->maxiter = min( 2 * A.num_cols, solver_par->maxiter );
// check if matrix A is square
if ( A.num_rows != A.num_cols ) {
//printf("Matrix A is not square.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
goto cleanup;
}
// |b|
nrmb = magma_dznrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_zscal( x->num_rows, MAGMA_Z_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// t = 0
// make t twice as large to contain both, dt and dr
ldd = magma_roundup( b.num_rows, 32 );
CHECK( magma_zvinit( &dt, Magma_DEV, ldd, 2, c_zero, queue ));
dt.num_rows = b.num_rows;
dt.num_cols = 1;
dt.nnz = dt.num_rows;
// redirect the dr.dval to the second part of dt
CHECK( magma_zvinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
magma_free( dr.dval );
dr.dval = dt.dval + ldd;
// r = b - A x
CHECK( magma_zresidualvec( A, b, *x, &dr, &nrmr, queue ));
// |r|
solver_par->init_res = nrmr;
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t)nrmr;
}
// check if initial is guess good enough
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
// P = randn(n, s)
// P = ortho(P)
//---------------------------------------
// P = 0.0
CHECK( magma_zvinit( &dP, Magma_CPU, A.num_cols, s, c_zero, queue ));
// P = randn(n, s)
distr = 3; // 1 = unif (0,1), 2 = unif (-1,1), 3 = normal (0,1)
dof = dP.num_rows * dP.num_cols;
lapackf77_zlarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_zmtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_zmfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_zqr(P1), QR factorization
CHECK( magma_zqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_dznrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_zdscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_zmtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_zmfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_zvinit( &hskp, Magma_CPU, 4, 1, c_zero, queue ));
CHECK( magma_zvinit( &dskp, Magma_DEV, 4, 1, c_zero, queue ));
CHECK( magma_zvinit( &halpha, Magma_CPU, s, 1, c_zero, queue ));
CHECK( magma_zvinit( &dalpha, Magma_DEV, s, 1, c_zero, queue ));
CHECK( magma_zvinit( &hbeta, Magma_CPU, s, 1, c_zero, queue ));
CHECK( magma_zvinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// workspace for merged dot product
CHECK( magma_zmalloc( &d1, max(2, s) * b.num_rows ));
CHECK( magma_zmalloc( &d2, max(2, s) * b.num_rows ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_zmtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// tt = 0
// make tt twice as large to contain both, dtt and drs
ldd = magma_roundup( b.num_rows, 32 );
CHECK( magma_zvinit( &dtt, Magma_DEV, ldd, 2, c_zero, queue ));
dtt.num_rows = dr.num_rows;
dtt.num_cols = 1;
dtt.nnz = dtt.num_rows;
// redirect the drs.dval to the second part of dtt
CHECK( magma_zvinit( &drs, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
magma_free( drs.dval );
drs.dval = dtt.dval + ldd;
// set smoothing residual vector
magma_zcopyvector( dr.num_rows, dr.dval, 1, drs.dval, 1, queue );
}
// G(n,s) = 0
if ( s > 1 ) {
ldd = magma_roundup( A.num_rows, 32 );
CHECK( magma_zvinit( &dG, Magma_DEV, ldd, s, c_zero, queue ));
dG.num_rows = A.num_rows;
} else {
CHECK( magma_zvinit( &dG, Magma_DEV, A.num_rows, s, c_zero, queue ));
}
// dGcol represents a single column of dG, array pointer is set inside loop
CHECK( magma_zvinit( &dGcol, Magma_DEV, dG.num_rows, 1, c_zero, queue ));
magma_free( dGcol.dval );
// U(n,s) = 0
if ( s > 1 ) {
ldd = magma_roundup( A.num_cols, 32 );
CHECK( magma_zvinit( &dU, Magma_DEV, ldd, s, c_zero, queue ));
dU.num_rows = A.num_cols;
} else {
CHECK( magma_zvinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
}
// M(s,s) = I
CHECK( magma_zvinit( &dM, Magma_DEV, s, s, c_zero, queue ));
CHECK( magma_zvinit( &hMdiag, Magma_CPU, s, 1, c_zero, queue ));
magmablas_zlaset( MagmaFull, dM.num_rows, dM.num_cols, c_zero, c_one, dM.dval, dM.ld, queue );
// f = 0
CHECK( magma_zvinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// c = 0
CHECK( magma_zvinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = 0
CHECK( magma_zvinit( &dv, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// lu = 0
CHECK( magma_zvinit( &dlu, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
om = MAGMA_Z_ONE;
innerflag = 0;
// start iteration
do
{
solver_par->numiter++;
// new RHS for small systems
// f = P' r
magma_zgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queue );
// shadow space loop
for ( k = 0; k < s; ++k ) {
sk = s - k;
// c(k:s) = M(k:s,k:s) \ f(k:s)
magma_zcopyvector( sk, &df.dval[k], 1, &dc.dval[k], 1, queue );
magma_ztrsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, sk, &dM.dval[k*dM.ld+k], dM.ld, &dc.dval[k], 1, queue );
// v = r - G(:,k:s) c(k:s)
magma_zcopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
magmablas_zgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, &dG.dval[k*dG.ld], dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queue );
// preconditioning operation
// v = L \ v;
// v = U \ v;
CHECK( magma_z_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queue ));
CHECK( magma_z_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queue ));
// U(:,k) = om * v + U(:,k:s) c(k:s)
magmablas_zgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queue );
magma_zcopyvector( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queue );
// G(:,k) = A U(:,k)
dGcol.dval = dG.dval + k * dG.ld;
CHECK( magma_z_spmv( c_one, A, dv, c_zero, dGcol, queue ));
solver_par->spmv_count++;
// bi-orthogonalize the new basis vectors
for ( i = 0; i < k; ++i ) {
// alpha = P(:,i)' G(:,k)
halpha.val[i] = magma_zdotc( dP.num_rows, &dP.dval[i*dP.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// alpha = alpha / M(i,i)
halpha.val[i] = halpha.val[i] / hMdiag.val[i];
// G(:,k) = G(:,k) - alpha * G(:,i)
magma_zaxpy( dG.num_rows, -halpha.val[i], &dG.dval[i*dG.ld], 1, &dG.dval[k*dG.ld], 1, queue );
}
// non-first s iteration
if ( k > 0 ) {
// U update outside of loop using GEMV
// U(:,k) = U(:,k) - U(:,1:k) * alpha(1:k)
magma_zsetvector( k, halpha.val, 1, dalpha.dval, 1, queue );
magmablas_zgemv( MagmaNoTrans, dU.num_rows, k, c_n_one, dU.dval, dU.ld, dalpha.dval, 1, c_one, &dU.dval[k*dU.ld], 1, queue );
}
// new column of M = P'G, first k-1 entries are zero
// M(k:s,k) = P(:,k:s)' G(:,k)
magma_zgemvmdot_shfl( dP.num_rows, sk, &dP.dval[k*dP.ld], &dG.dval[k*dG.ld], d1, d2, &dM.dval[k*dM.ld+k], queue );
magma_zgetvector( 1, &dM.dval[k*dM.ld+k], 1, &hMdiag.val[k], 1, queue );
// check M(k,k) == 0
if ( MAGMA_Z_EQUAL(hMdiag.val[k], MAGMA_Z_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
magma_zgetvector( 1, &df.dval[k], 1, &fk, 1, queue );
hbeta.val[k] = fk / hMdiag.val[k];
// check for nan
if ( magma_z_isnan( hbeta.val[k] ) || magma_z_isinf( hbeta.val[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
magma_zaxpy( dr.num_rows, -hbeta.val[k], &dG.dval[k*dG.ld], 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
nrmr = magma_dznrm2( dr.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// x = x + beta * U(:,k)
magma_zaxpy( x->num_rows, hbeta.val[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queue );
// smoothing operation
//---------------------------------------
// t = rs - r
magma_zidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queue );
// t't
// t'rs
CHECK( magma_zgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queue ));
magma_zgetvector( 2, &dskp.dval[2], 1, &hskp.val[2], 1, queue );
// gamma = (t' * rs) / (t' * t)
gamma = hskp.val[3] / hskp.val[2];
// rs = rs - gamma * (rs - r)
magma_zaxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_zidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queue );
// |rs|
nrmr = magma_dznrm2( drs.num_rows, drs.dval, 1, queue );
//---------------------------------------
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter) % solver_par->verbose == 0 ) {
solver_par->res_vec[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence or iteration limit
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
s = k + 1; // for the x-update outside the loop
innerflag = 2;
info = MAGMA_SUCCESS;
break;
}
// non-last s iteration
if ( (k + 1) < s ) {
// f(k+1:s) = f(k+1:s) - beta * M(k+1:s,k)
magma_zaxpy( sk-1, -hbeta.val[k], &dM.dval[k*dM.ld+(k+1)], 1, &df.dval[k+1], 1, queue );
}
}
// smoothing disabled
if ( smoothing <= 0 && innerflag != 1 ) {
// update solution approximation x
// x = x + U(:,1:s) * beta(1:s)
magma_zsetvector( s, hbeta.val, 1, dbeta.dval, 1, queue );
magmablas_zgemv( MagmaNoTrans, dU.num_rows, s, c_one, dU.dval, dU.ld, dbeta.dval, 1, c_one, x->dval, 1, queue );
}
// check convergence or iteration limit or invalid result of inner loop
if ( innerflag > 0 ) {
break;
}
// v = r
magma_zcopy( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
// preconditioning operation
// v = L \ v;
// v = U \ v;
CHECK( magma_z_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queue ));
CHECK( magma_z_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queue ));
// t = A v
CHECK( magma_z_spmv( c_one, A, dv, c_zero, dt, queue ));
solver_par->spmv_count++;
// computation of a new omega
//---------------------------------------
// t't
// t'r
CHECK( magma_zgemvmdot_shfl( dt.ld, 2, dt.dval, dt.dval, d1, d2, dskp.dval, queue ));
magma_zgetvector( 2, dskp.dval, 1, hskp.val, 1, queue );
// |t|
nrmt = magma_dsqrt( MAGMA_Z_REAL(hskp.val[0]) );
// rho = abs((t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_Z_REAL(hskp.val[1]) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = hskp.val[1] / hskp.val[0];
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_Z_EQUAL(om, MAGMA_Z_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// update approximation vector
// x = x + om * v
magma_zaxpy( x->num_rows, om, dv.dval, 1, x->dval, 1, queue );
// update residual vector
// r = r - om * t
magma_zaxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// residual norm
nrmr = magma_dznrm2( dr.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
magma_zidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queue );
// t't
// t'rs
CHECK( magma_zgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queue ));
magma_zgetvector( 2, &dskp.dval[2], 1, &hskp.val[2], 1, queue );
// gamma = (t' * rs) / (t' * t)
gamma = hskp.val[3] / hskp.val[2];
// rs = rs - gamma * (rs - r)
magma_zaxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_zidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queue );
// |rs|
nrmr = magma_dznrm2( drs.num_rows, drs.dval, 1, queue );
//---------------------------------------
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter) % solver_par->verbose == 0 ) {
solver_par->res_vec[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
break;
}
}
while ( solver_par->numiter + 1 <= solver_par->maxiter );
// smoothing enabled
if ( smoothing > 0 ) {
// x = xs
magma_zcopyvector( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_zcopyvector( dr.num_rows, drs.dval, 1, dr.dval, 1, queue );
}
// get last iteration timing
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t)tempo2 - tempo1;
//--------------STOP TIME----------------
// get final stats
solver_par->iter_res = nrmr;
CHECK( magma_zresidualvec( A, b, *x, &dr, &residual, queue ));
solver_par->final_res = residual;
// set solver conclusion
if ( info != MAGMA_SUCCESS && info != MAGMA_DIVERGENCE ) {
if ( solver_par->init_res > solver_par->final_res ) {
info = MAGMA_SLOW_CONVERGENCE;
}
}
cleanup:
// free resources
// smoothing enabled
if ( smoothing > 0 ) {
drs.dval = NULL; // needed because its pointer is redirected to dtt
magma_zmfree( &dxs, queue );
magma_zmfree( &drs, queue );
magma_zmfree( &dtt, queue );
}
dr.dval = NULL; // needed because its pointer is redirected to dt
dGcol.dval = NULL; // needed because its pointer is redirected to dG
magma_zmfree( &dr, queue );
magma_zmfree( &dP, queue );
magma_zmfree( &dP1, queue );
magma_zmfree( &dG, queue );
magma_zmfree( &dGcol, queue );
magma_zmfree( &dU, queue );
magma_zmfree( &dM, queue );
magma_zmfree( &hMdiag, queue );
magma_zmfree( &df, queue );
magma_zmfree( &dt, queue );
magma_zmfree( &dc, queue );
magma_zmfree( &dv, queue );
magma_zmfree( &dlu, queue );
magma_zmfree( &dskp, queue );
magma_zmfree( &dalpha, queue );
magma_zmfree( &dbeta, queue );
magma_zmfree( &hskp, queue );
magma_zmfree( &halpha, queue );
magma_zmfree( &hbeta, queue );
magma_free( d1 );
magma_free( d2 );
solver_par->info = info;
return info;
/* magma_zpidr_merge */
}
extern "C" magma_int_t
magma_zgeqrs3_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *tau, magmaDoubleComplex *dT,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *hwork, magma_int_t lwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by ZGEQRF3_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) COMPLEX_16 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
ZGEQRF3_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) COMPLEX_16 array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_ZGEQRF_GPU.
DB (input/output) COMPLEX_16 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) COMPLEX_16 array that is the output (the 6th argument)
of magma_zgeqrf_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
matrices 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) COMPLEX_16 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_zgeqrf_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+(a_2)*(ldda) + (a_1))
#define d_ref(a_1) (dT+(lddwork+(a_1))*nb)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t k, lddwork;
magma_int_t nb = magma_get_zgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_Z_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;
}
lddwork= k;
/* B := Q' * B */
magma_zunmqr_gpu( MagmaLeft, MagmaConjTrans,
m, nrhs, n,
a_ref(0,0), ldda, tau,
dB, lddb, hwork, lwork, dT, nb, info );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:)
1. Move the block diagonal submatrices from d_ref to R
2. Solve
3. Restore the data format moving data from R back to d_ref
*/
magmablas_zswapdblk(k, nb, a_ref(0,0), ldda, 1, d_ref(0), nb, 0);
if ( nrhs == 1 ) {
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, a_ref(0,0), ldda, dB, 1);
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, nrhs, c_one, a_ref(0,0), ldda, dB, lddb);
}
magmablas_zswapdblk(k, nb, d_ref(0), nb, 0, a_ref(0,0), ldda, 1);
return *info;
}
extern "C" magma_err_t
magma_zgetrs_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_int_t *ipiv,
magmaDoubleComplex_ptr dB, size_t dB_offset, magma_int_t lddb,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMagma (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
April 2012
Purpose
=======
Solves a system of linear equations
A * X = B or A' * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The order of the matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
A (input) COMPLEX_16 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
B (input/output) COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
HWORK (workspace) COMPLEX_16 array, dimension N*NRHS
===================================================================== */
magmaDoubleComplex z_one = MAGMA_Z_MAKE( 1.0, 0.0 );
magmaDoubleComplex *work = NULL;
magma_trans_t trans_ = trans;
long int notran = lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaNoTrans));
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(! lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaTrans))) &&
(! lapackf77_lsame(lapack_const(trans_), lapack_const(MagmaConjTrans))) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (n == 0 || nrhs == 0) {
return *info;
}
work=(magmaDoubleComplex*)malloc( n*nrhs*sizeof(magmaDoubleComplex) );
if ( !work ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
i1 = 1;
i2 = n;
if (notran) {
inc = 1;
/* Solve A * X = B. */
magma_zgetmatrix( n, nrhs, dB, dB_offset, lddb, work, 0, n, queue );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, 0, n, dB, dB_offset, lddb, queue );
if ( nrhs == 1) {
chk(magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue));
chk(magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue));
} else {
chk(magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
chk(magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
}
} else {
inc = -1;
/* Solve A' * X = B. */
if ( nrhs == 1) {
chk(magma_ztrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue ));
chk(magma_ztrsv(MagmaLower, trans, MagmaUnit, n, dA, dA_offset, ldda, dB, dB_offset, 1, queue ));
} else {
chk(magma_ztrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
chk(magma_ztrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, z_one, dA, dA_offset, ldda, dB, dB_offset, lddb, queue));
}
magma_zgetmatrix( n, nrhs, dB, dB_offset, lddb, work, 0, n, queue );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, 0, n, dB, dB_offset, lddb, queue );
}
free(work);
return *info;
}
/**
Purpose
-------
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by ZGEQRF3_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 COMPLEX_16 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
ZGEQRF3_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 COMPLEX_16 array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_ZGEQRF_GPU.
@param[in,out]
dB COMPLEX_16 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 COMPLEX_16 array that is the output (the 6th argument)
of magma_zgeqrf_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
matrices 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) COMPLEX_16 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_zgeqrf_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_zgels_comp
********************************************************************/
extern "C" magma_int_t
magma_zgeqrs3_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *tau, magmaDoubleComplex *dT,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *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)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t k, lddwork;
magma_int_t nb = magma_get_zgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_Z_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;
}
lddwork= k;
/* B := Q' * B */
magma_zunmqr_gpu( MagmaLeft, MagmaConjTrans,
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,:)
1. Move the block diagonal submatrices from dT to R
2. Solve
3. Restore the data format moving data from R back to dT
*/
magmablas_zswapdblk(k, nb, dA(0,0), ldda, 1, dT(0), nb, 0);
if ( nrhs == 1 ) {
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, dA(0,0), ldda, dB, 1);
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, nrhs, c_one, dA(0,0), ldda, dB, lddb);
}
magmablas_zswapdblk(k, nb, dT(0), nb, 0, dA(0,0), ldda, 1);
return *info;
}
extern "C" magma_int_t
magma_zpotrs_gpu(char uplo, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dB, magma_int_t lddb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZPOTRS solves a system of linear equations A*X = B with a Hermitian
positive definite matrix A using the Cholesky factorization
A = U**H*U or A = L*L**H computed by ZPOTRF.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
dA (input) COMPLEX_16 array on the GPU, dimension (LDDA,N)
The triangular factor U or L from the Cholesky factorization
A = U**H*U or A = L*L**H, as computed by ZPOTRF.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
dB (input/output) COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
LDDB (input) INTEGER
The leading dimension of the array B. LDDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
magmaDoubleComplex c_one = MAGMA_Z_ONE;
*info = 0 ;
if( (uplo != 'U') && (uplo != 'u') && (uplo != 'L') && (uplo != 'l') )
*info = -1;
if( n < 0 )
*info = -2;
if( nrhs < 0)
*info = -3;
if ( ldda < max(1, n) )
*info = -5;
if ( lddb < max(1, n) )
*info = -7;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if ( (n == 0) || (nrhs == 0) ) {
return *info;
}
if( (uplo=='U') || (uplo=='u') ){
if ( nrhs == 1) {
magma_ztrsv(MagmaUpper, MagmaConjTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
}
}
else{
if ( nrhs == 1) {
magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaLower, MagmaConjTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb);
}
}
return *info;
}
extern "C" magma_int_t
magma_zgetrs_gpu(char trans, magma_int_t n, magma_int_t nrhs,
cuDoubleComplex *dA, magma_int_t ldda,
magma_int_t *ipiv,
cuDoubleComplex *dB, magma_int_t lddb,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
Solves a system of linear equations
A * X = B or A' * X = B
with a general N-by-N matrix A using the LU factorization computed by ZGETRF_GPU.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The order of the matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
A (input) COMPLEX_16 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by ZGETRF_GPU.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
IPIV (input) INTEGER array, dimension (N)
The pivot indices from ZGETRF; for 1<=i<=N, row i of the
matrix was interchanged with row IPIV(i).
B (input/output) COMPLEX_16 array on the GPU, dimension (LDB,NRHS)
On entry, the right hand side matrix B.
On exit, the solution matrix X.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
HWORK (workspace) COMPLEX_16 array, dimension N*NRHS
===================================================================== */
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex *work = NULL;
char trans_[2] = {trans, 0};
int notran = lapackf77_lsame(trans_, "N");
magma_int_t i1, i2, inc;
*info = 0;
if ( (! notran) &&
(! lapackf77_lsame(trans_, "T")) &&
(! lapackf77_lsame(trans_, "C")) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (n == 0 || nrhs == 0) {
return *info;
}
magma_zmalloc_cpu( &work, n * nrhs );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
i1 = 1;
i2 = n;
if (notran) {
inc = 1;
/* Solve A * X = B. */
magma_zgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, n, dB, lddb );
if ( nrhs == 1) {
magma_ztrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
} else {
inc = -1;
/* Solve A' * X = B. */
if ( nrhs == 1) {
magma_ztrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ztrsv(MagmaLower, trans, MagmaUnit, n, dA, ldda, dB, 1 );
} else {
magma_ztrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ztrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
magma_zgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_zlaswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_zsetmatrix( n, nrhs, work, n, dB, lddb );
}
magma_free_cpu(work);
return *info;
}
