magma_int_t
magma_vector_zlag2c( magma_z_vector x, magma_c_vector *y )
{
magma_int_t info;
if( x.memory_location == Magma_DEV){
y->memory_location = x.memory_location;
y->num_rows = x.num_rows;
y->nnz = x.nnz;
magma_cmalloc( &y->val, x.num_rows );
magmablas_zlag2c_sparse( x.num_rows, 1, x.val, x.num_rows, y->val,
x.num_rows, &info );
return MAGMA_SUCCESS;
}
else if( x.memory_location == Magma_CPU ){
y->memory_location = x.memory_location;
y->num_rows = x.num_rows;
y->nnz = x.nnz;
magma_cmalloc_cpu( &y->val, x.num_rows );
magma_int_t one= 1;
magma_int_t info;
lapackf77_zlag2c( &x.num_rows, &one,
x.val, &x.num_rows,
y->val, &x.num_rows, &info);
return MAGMA_SUCCESS;
}
else
return MAGMA_ERR_NOT_SUPPORTED;
}
extern "C" void
magma_clarfxsym(
magma_int_t N,
magmaFloatComplex *A, magma_int_t LDA,
magmaFloatComplex *V, magmaFloatComplex *TAU)
{
magma_int_t IONE=1;
magmaFloatComplex dtmp;
magmaFloatComplex Z_ZERO = MAGMA_C_ZERO;
//magmaFloatComplex Z_ONE = MAGMA_C_ONE;
magmaFloatComplex Z_MONE = MAGMA_C_NEG_ONE;
magmaFloatComplex Z_HALF = MAGMA_C_HALF;
//magmaFloatComplex WORK[N];
magmaFloatComplex *WORK;
magma_cmalloc_cpu( &WORK, N );
/* apply left and right on A(st:ed,st:ed)*/
//magma_clarfxsym(len,A(st,st),LDX,V(st),TAU(st));
/* X = AVtau */
blasf77_chemv("L",&N, TAU, A, &LDA, V, &IONE, &Z_ZERO, WORK, &IONE);
/* je calcul dtmp= X'*V */
dtmp = magma_cblas_cdotc(N, WORK, IONE, V, IONE);
/* je calcul 1/2 X'*V*t = 1/2*dtmp*tau */
dtmp = -dtmp * Z_HALF * (*TAU);
/* je calcul W=X-1/2VX'Vt = X - dtmp*V */
/*
for (j = 0; j < N; j++)
WORK[j] = WORK[j] + (dtmp*V[j]); */
blasf77_caxpy(&N, &dtmp, V, &IONE, WORK, &IONE);
/* performs the symmetric rank 2 operation A := alpha*x*y' + alpha*y*x' + A */
blasf77_cher2("L",&N,&Z_MONE,WORK,&IONE,V,&IONE,A,&LDA);
magma_free_cpu(WORK);
}
extern "C" void
magma_ctrdtype3cbHLsym_withQ(
magma_int_t N, magma_int_t NB,
magmaFloatComplex *A, magma_int_t LDA,
magmaFloatComplex *V, magmaFloatComplex *TAU,
magma_int_t st, magma_int_t ed, magma_int_t sweep, magma_int_t Vblksiz)
{
//magma_int_t J1, J2, J3, i, j;
magma_int_t len, LDX;
//magma_int_t IONE=1;
magma_int_t blkid, vpos, taupos, tpos;
//magmaFloatComplex conjtmp;
magmaFloatComplex *WORK;
magma_cmalloc_cpu( &WORK, N );
findVTpos(N,NB,Vblksiz,sweep-1,st-1, &vpos, &taupos, &tpos, &blkid);
LDX = LDA-1;
len = ed-st+1;
/* apply left and right on A(st:ed,st:ed)*/
magma_clarfxsym(len,A(st,st),LDX,V(vpos),TAU(taupos));
//conjtmp = MAGMA_C_CONJ(*TAU(taupos));
//lapackf77_clarfy("L", &len, V(vpos), &IONE, &(MAGMA_C_CONJ(*TAU(taupos))), A(st,st), &LDX, WORK);
magma_free_cpu(WORK);
}
extern "C" void
magma_ctrdtype1cbHLsym_withQ(
magma_int_t N, magma_int_t NB,
magmaFloatComplex *A, magma_int_t LDA,
magmaFloatComplex *V, magmaFloatComplex *TAU,
magma_int_t st, magma_int_t ed, magma_int_t sweep, magma_int_t Vblksiz)
{
//magma_int_t J1, J2, J3, i, j;
magma_int_t len, LDX;
magma_int_t IONE=1;
magma_int_t blkid, vpos, taupos, tpos;
//magmaFloatComplex conjtmp;
magmaFloatComplex Z_ONE = MAGMA_C_ONE;
magmaFloatComplex *WORK;
magma_cmalloc_cpu( &WORK, N );
findVTpos(N,NB,Vblksiz,sweep-1,st-1, &vpos, &taupos, &tpos, &blkid);
//printf("voici vpos %d taupos %d tpos %d blkid %d \n", vpos, taupos, tpos, blkid);
LDX = LDA-1;
len = ed-st+1;
*V(vpos) = Z_ONE;
memcpy(V(vpos+1), A(st+1, st-1), (len-1)*sizeof(magmaFloatComplex));
memset(A(st+1, st-1), 0, (len-1)*sizeof(magmaFloatComplex));
/* Eliminate the col at st-1 */
lapackf77_clarfg( &len, A(st, st-1), V(vpos+1), &IONE, TAU(taupos) );
/* apply left and right on A(st:ed,st:ed)*/
magma_clarfxsym(len,A(st,st),LDX,V(vpos),TAU(taupos));
//conjtmp = MAGMA_C_CONJ(*TAU(taupos));
//lapackf77_clarfy("L", &len, V(vpos), &IONE, &conjtmp, A(st,st), &LDX, WORK); //&(MAGMA_C_CONJ(*TAU(taupos)))
magma_free_cpu(WORK);
}
void magma_cprint_gpu(
magma_int_t m, magma_int_t n,
const magmaFloatComplex *dA, magma_int_t ldda )
{
magma_int_t info = 0;
if ( m < 0 )
info = -1;
else if ( n < 0 )
info = -2;
else if ( ldda < max(1,m) )
info = -4;
if (info != 0) {
magma_xerbla( __func__, -(info) );
return; //info;
}
magma_int_t lda = m;
magmaFloatComplex* A;
magma_cmalloc_cpu( &A, lda*n );
magma_cgetmatrix( m, n, dA, ldda, A, lda );
magma_cprint( m, n, A, lda );
magma_free_cpu( A );
}
magma_int_t magma_cnan_inf_gpu(
magma_uplo_t uplo, magma_int_t m, magma_int_t n,
const magmaFloatComplex *dA, magma_int_t ldda,
magma_int_t *cnt_nan,
magma_int_t *cnt_inf )
{
magma_int_t info = 0;
if ( uplo != MagmaLower && uplo != MagmaUpper && uplo != MagmaFull )
info = -1;
else if ( m < 0 )
info = -2;
else if ( n < 0 )
info = -3;
else if ( magma_is_devptr( dA ) == 0 )
info = -4;
else if ( ldda < max(1,m) )
info = -5;
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
magma_int_t lda = m;
magmaFloatComplex* A;
magma_cmalloc_cpu( &A, lda*n );
magma_cgetmatrix( m, n, dA, ldda, A, lda );
magma_int_t cnt = magma_cnan_inf( uplo, m, n, A, lda, cnt_nan, cnt_inf );
magma_free_cpu( A );
return cnt;
}
void magma_cprint_gpu( magma_int_t m, magma_int_t n, const magmaFloatComplex *dA, magma_int_t ldda )
{
if ( magma_is_devptr( dA ) == 0 ) {
fprintf( stderr, "ERROR: cprint_gpu called with host pointer.\n" );
exit(1);
}
magma_int_t lda = m;
magmaFloatComplex* A;
magma_cmalloc_cpu( &A, lda*n );
magma_cgetmatrix( m, n, dA, ldda, A, lda );
magma_cprint( m, n, A, lda );
magma_free_cpu( A );
}
//////////////////////////////////////////////////////////////
// CSTEDC Divide and Conquer for tridiag
//////////////////////////////////////////////////////////////
extern "C" void magma_cstedc_withZ(magma_vec_t JOBZ, magma_int_t N, float *D, float * E, magmaFloatComplex *Z, magma_int_t LDZ)
{
magmaFloatComplex *WORK;
float *RWORK;
magma_int_t *IWORK;
magma_int_t LWORK, LIWORK, LRWORK;
magma_int_t INFO;
// use log() as log2() is not defined everywhere (e.g., Windows)
const float log_2 = 0.6931471805599453;
if (JOBZ == MagmaVec) {
LWORK = N*N;
LRWORK = 1 + 3*N + 3*N*((magma_int_t)(log( (float)N )/log_2) + 1) + 4*N*N + 256*N;
LIWORK = 6 + 6*N + 6*N*((magma_int_t)(log( (float)N )/log_2) + 1) + 256*N;
} else if (JOBZ == MagmaIVec) {
LWORK = N;
LRWORK = 2*N*N + 4*N + 1 + 256*N;
LIWORK = 256*N;
} else if (JOBZ == MagmaNoVec) {
LWORK = N;
LRWORK = 256*N + 1;
LIWORK = 256*N;
} else {
printf("ERROR JOBZ %c\n", JOBZ);
exit(-1);
}
magma_smalloc_cpu( &RWORK, LRWORK );
magma_cmalloc_cpu( &WORK, LWORK );
magma_imalloc_cpu( &IWORK, LIWORK );
lapackf77_cstedc( lapack_vec_const(JOBZ), &N, D, E, Z, &LDZ, WORK, &LWORK, RWORK, &LRWORK, IWORK, &LIWORK, &INFO);
if (INFO != 0) {
printf("=================================================\n");
printf("CSTEDC ERROR OCCURED. HERE IS INFO %d \n ", (int) INFO);
printf("=================================================\n");
//assert(INFO == 0);
}
magma_free_cpu( IWORK );
magma_free_cpu( WORK );
magma_free_cpu( RWORK );
}
extern "C" void
magma_ctrdtype2cbHLsym_withQ(
magma_int_t N, magma_int_t NB,
magmaFloatComplex *A, magma_int_t LDA,
magmaFloatComplex *V, magmaFloatComplex *TAU,
magma_int_t st, magma_int_t ed, magma_int_t sweep, magma_int_t Vblksiz)
{
magma_int_t J1, J2, len, lem, LDX;
//magma_int_t i, j;
magma_int_t IONE=1;
magma_int_t blkid, vpos, taupos, tpos;
magmaFloatComplex conjtmp;
magmaFloatComplex Z_ONE = MAGMA_C_ONE;
//magmaFloatComplex WORK[NB];
magmaFloatComplex *WORK;
magma_cmalloc_cpu( &WORK, NB );
findVTpos(N,NB,Vblksiz,sweep-1,st-1, &vpos, &taupos, &tpos, &blkid);
LDX = LDA-1;
J1 = ed+1;
J2 = min(ed+NB,N);
len = ed-st+1;
lem = J2-J1+1;
if (lem > 0) {
/* apply remaining right commming from the top block */
lapackf77_clarfx("R", &lem, &len, V(vpos), TAU(taupos), A(J1, st), &LDX, WORK);
}
if (lem > 1) {
findVTpos(N,NB,Vblksiz,sweep-1,J1-1, &vpos, &taupos, &tpos, &blkid);
/* remove the first column of the created bulge */
*V(vpos) = Z_ONE;
memcpy(V(vpos+1), A(J1+1, st), (lem-1)*sizeof(magmaFloatComplex));
memset(A(J1+1, st),0,(lem-1)*sizeof(magmaFloatComplex));
/* Eliminate the col at st */
lapackf77_clarfg( &lem, A(J1, st), V(vpos+1), &IONE, TAU(taupos) );
/* apply left on A(J1:J2,st+1:ed) */
len = len-1; /* because we start at col st+1 instead of st. col st is the col that has been revomved; */
conjtmp = MAGMA_C_CONJ(*TAU(taupos));
lapackf77_clarfx("L", &lem, &len, V(vpos), &conjtmp, A(J1, st+1), &LDX, WORK);
}
magma_free_cpu(WORK);
}
magma_int_t magma_cnan_inf_gpu(
magma_uplo_t uplo, magma_int_t m, magma_int_t n,
magmaFloatComplex_const_ptr dA, magma_int_t ldda,
magma_int_t *cnt_nan,
magma_int_t *cnt_inf )
{
magma_int_t info = 0;
if ( uplo != MagmaLower && uplo != MagmaUpper && uplo != MagmaFull )
info = -1;
else if ( m < 0 )
info = -2;
else if ( n < 0 )
info = -3;
else if ( ldda < max(1,m) )
info = -5;
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
magma_int_t lda = m;
magmaFloatComplex* A;
magma_cmalloc_cpu( &A, lda*n );
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
magma_cgetmatrix( m, n, dA, ldda, A, lda, queue );
magma_queue_destroy( queue );
magma_int_t cnt = magma_cnan_inf( uplo, m, n, A, lda, cnt_nan, cnt_inf );
magma_free_cpu( A );
return cnt;
}
extern "C" magma_int_t
magma_cgetrf_gpu(magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, 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.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
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. N >= 0.
A (input/output) COMPLEX array on the GPU, dimension (LDDA,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.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
> 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.
===================================================================== */
#define dAT(i,j) (dAT + (i)*nb*lddat + (j)*nb)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, rows, cols, s, lddat, lddwork;
magmaFloatComplex *dAT, *dAP, *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < 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;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
lddat = maxn;
lddwork = maxm;
dAT = dA;
if (MAGMA_SUCCESS != magma_cmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if ( m == n ) {
lddat = ldda;
magmablas_ctranspose_inplace( m, dAT, ldda );
}
else {
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ctranspose2( dAT, lddat, dA, ldda, m, n );
}
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
magma_free( dAP );
if ( ! (m == n))
magma_free( dAT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
cudaStream_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else
stream[1] = current_stream;
for( i=0; i<s; i++ )
{
// download i-th panel
cols = maxm - i*nb;
//magmablas_ctranspose( dAP, cols, dAT(i,i), lddat, nb, cols );
magmablas_ctranspose2( dAP, cols, dAT(i,i), lddat, nb, m-i*nb );
// make sure that that the transpose has completed
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dAP, cols, work, lddwork,
stream[0]);
if ( i>0 ){
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), lddat,
dAT(i-1,i+1), lddat );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), lddat,
dAT(i, i-1), lddat,
c_one, dAT(i, i+1), lddat );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lddwork, ipiv+i*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dAP, maxm,
stream[0]);
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
//magmablas_ctranspose(dAT(i,i), lddat, dAP, maxm, cols, nb);
magmablas_ctranspose2(dAT(i,i), lddat, dAP, maxm, m-i*nb, nb);
// do the small non-parallel computations (next panel update)
if ( s > (i+1) ) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose2( dAP, maxm, dAT(s,s), lddat, nb0, rows);
magma_cgetmatrix( rows, nb0, dAP, maxm, work, lddwork );
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lddwork, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb0, s*nb );
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dAP, maxm );
magmablas_ctranspose2( dAT(s,s), lddat, dAP, maxm, rows, nb0);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
if ( m == n ) {
magmablas_ctranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ctranspose2( dA, ldda, dAT, lddat, n, m );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* End of MAGMA_CGETRF_GPU */
extern "C" magma_int_t
magma_cungqr(magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *dT, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
DT (input) COMPLEX array on the GPU device.
DT contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu.
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk; //, iinfo;
magma_int_t lddwork;
magmaFloatComplex *dA, *dV, *dW;
magmaFloatComplex *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate GPU work space
// ldda*n for matrix dA
// ldda*nb for dV
// lddwork*nb for dW larfb workspace
ldda = ((m + 31) / 32) * 32;
lddwork = ((n + 31) / 32) * 32;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, ldda*n + ldda*nb + lddwork*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dA + ldda*n;
dW = dA + ldda*n + ldda*nb;
// Allocate CPU work space
lwork = (n+m+nb) * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
magma_free( dA );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmaFloatComplex *V = work + (n+nb)*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
/*
// Replacing this with the following 4 routines works but cungqr is slow for
// k smaller than the cungqr's blocking size (new version can be up to 60x faster)
lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_clacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_claset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_clarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
V, &m_kk, work, &k_kk, A(kk, kk), &lda, work+k_kk*k_kk, &n_kk );
if (kk > 0) {
magma_csetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaUpperLower, kk, n - kk, dA(0, kk), ldda );
}
}
if (kk > 0) {
// Use blocked code
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min(nb, k - i);
// Send current panel to the GPU
mi = m - i;
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
// set panel to identity
magmablas_claset( MagmaUpperLower, i, ib, dA(0, i), ldda );
magmablas_claset_identity( mi, ib, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
// copy result back to CPU
magma_cgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda);
}
magmablasSetKernelStream( NULL );
magma_queue_destroy( stream );
magma_free( dA );
magma_free_cpu( work );
return *info;
} /* magma_cungqr */
extern "C" magma_int_t
magma_cungqr_m(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *T, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
T (input) COMPLEX array, dimension (NB, min(M,N)).
T contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu (except stored on the CPU, not the GPU).
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in T.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(d,i,j) (dA[d] + (i) + (j)*ldda)
#define dT(d,i,j) (dT[d] + (i) + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldwork;
magma_int_t i, ib, ki, kk, iinfo;
magmaFloatComplex *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
magma_int_t di, dn;
int dpanel;
int ngpu = magma_num_gpus();
int doriginal;
magma_getdevice( &doriginal );
// Allocate memory on GPUs for A and workspaces
magma_int_t ldda = ((m + 31) / 32) * 32;
magma_int_t lddwork = ((n + 31) / 32) * 32;
magma_int_t min_lblocks = (n / nb) / ngpu; // min. blocks per gpu
magma_int_t last_dev = (n / nb) % ngpu; // device with last block
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magmaFloatComplex *dA[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dT[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dV[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dW[ MagmaMaxGPUs ] = { NULL };
magma_queue_t stream[ MagmaMaxGPUs ] = { NULL };
for( int d = 0; d < ngpu; ++d ) {
// example with n = 75, nb = 10, ngpu = 3
// min_lblocks = 2
// last_dev = 1
// gpu 0: 2 blocks, cols: 0- 9, 30-39, 60-69
// gpu 1: 1+ blocks, cols: 10-19, 40-49, 70-74 (partial)
// gpu 2: 1 block , cols: 20-29, 50-59
magma_setdevice( d );
nlocal[d] = min_lblocks*nb;
if ( d < last_dev ) {
nlocal[d] += nb;
}
else if ( d == last_dev ) {
nlocal[d] += (n % nb);
}
ldwork = nlocal[d]*ldda // dA
+ nb*m // dT
+ nb*ldda // dV
+ nb*lddwork; // dW
if ( MAGMA_SUCCESS != magma_cmalloc( &dA[d], ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto CLEANUP;
}
dT[d] = dA[d] + nlocal[d]*ldda;
dV[d] = dT[d] + nb*m;
dW[d] = dV[d] + nb*ldda;
magma_queue_create( &stream[d] );
}
trace_init( 1, ngpu, 1, stream );
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for cungqr workspace
lwork = n * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
goto CLEANUP;
}
// Use unblocked code for the last or only block.
if (kk < n) {
trace_cpu_start( 0, "ungqr", "ungqr last block" );
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
dpanel = (kk / nb) % ngpu;
di = ((kk / nb) / ngpu) * nb;
magma_setdevice( dpanel );
lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
magma_csetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(dpanel, kk, di), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaUpperLower, kk, n - kk, dA(dpanel, 0, di), ldda );
trace_cpu_end( 0 );
}
if (kk > 0) {
// Use blocked code
// send T to all GPUs
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set T" );
magma_csetmatrix_async( nb, min(m,n), T, nb, dT[d], nb, stream[d] );
trace_gpu_end( d, 0 );
}
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
for( i = ki; i >= 0; i -= nb ) {
ib = min(nb, k - i);
mi = m - i;
dpanel = (i / nb) % ngpu;
di = ((i / nb) / ngpu) * nb;
// Send current panel to the GPUs
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set V" );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV[d], ldda, stream[d] );
trace_gpu_end( d, 0 );
}
// set panel to identity
magma_setdevice( dpanel );
magmablasSetKernelStream( stream[dpanel] );
trace_gpu_start( dpanel, 0, "laset", "laset" );
magmablas_claset( MagmaUpperLower, i, ib, dA(dpanel, 0, di), ldda );
magmablas_claset_identity( mi, ib, dA(dpanel, i, di), ldda );
trace_gpu_end( dpanel, 0 );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( stream[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, i, n, &di, &dn );
trace_gpu_start( d, 0, "larfb", "larfb" );
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, dn-di, ib,
dV[d], ldda, dT(d,0,i), nb,
dA(d, i, di), ldda, dW[d], lddwork );
trace_gpu_end( d, 0 );
}
}
}
}
// copy result back to CPU
trace_cpu_start( 0, "get", "get A" );
magma_cgetmatrix_1D_col_bcyclic( m, n, dA, ldda, A, lda, ngpu, nb );
trace_cpu_end( 0 );
#ifdef TRACING
char name[80];
snprintf( name, sizeof(name), "cungqr-n%d-ngpu%d.svg", m, ngpu );
trace_finalize( name, "trace.css" );
#endif
CLEANUP:
for( int d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( NULL );
magma_free( dA[d] );
dA[d] = NULL;
if ( stream[d] != NULL ) {
magma_queue_destroy( stream[d] );
}
}
magma_free_cpu( work );
magma_setdevice( doriginal );
return *info;
} /* magma_cungqr */
extern "C" magma_int_t
magma_cungqr(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau, magmaFloatComplex_ptr dT, size_t dT_offset,
magma_int_t nb,
magma_queue_t queue,
magma_int_t *info )
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
DT (input) COMPLEX array on the GPU device.
DT contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu.
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define a_ref(i,j) ( a + (j)*lda + (i))
#define da_ref(i,j) da, (da_offset + (j)*ldda + (i))
#define t_ref(a_1) dT, (dT_offset + (a_1)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magma_int_t i__1, i__2, i__3;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork = min(m, n);
magmaFloatComplex *work;
magmaFloatComplex_ptr da, dwork;
size_t da_offset, dwork_offset;
magma_event_t event = NULL;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0)
return *info;
/* Allocate GPU work space */
ldda = ((m+31)/32)*32;
lddwork = ((lddwork+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &da, ((n)*ldda + nb*lddwork ) )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
da_offset = 0;
dwork = da;
dwork_offset = da_offset + (n)*ldda;
/* Allocate CPU work space */
lwork = n * nb;
magma_cmalloc_cpu( &work, lwork );
if( work == NULL ) {
magma_free( da );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if ( (nb > 1) && (nb < k) )
{
/* Use blocked code after the last block.
The first kk columns are handled by the block method. */
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
/* Set A(1:kk,kk+1:n) to zero. */
magmablas_claset(MagmaFull, kk, n-kk, c_zero, c_zero, da_ref(0,kk), ldda, queue);
}
else
kk = 0;
/* Use unblocked code for the last or only block. */
if (kk < n)
{
i__1 = m - kk;
i__2 = n - kk;
i__3 = k - kk;
lapackf77_cungqr(&i__1, &i__2, &i__3,
a_ref(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo);
magma_csetmatrix(i__1, i__2, a_ref(kk, kk), lda, da_ref(kk, kk), ldda, queue);
}
if (kk > 0)
{
/* Use blocked code */
for (i = ki; i >= 0; i-=nb)
{
ib = min(nb, k - i);
/* Send the current panel to the GPU */
i__2 = m - i;
cpanel_to_q(MagmaUpper, ib, a_ref(i,i), lda, work);
magma_csetmatrix(i__2, ib, a_ref(i, i), lda, da_ref(i, i), ldda, queue);
if (i + ib < n)
{
/* Apply H to A(i:m,i+ib:n) from the left */
i__3 = n - i - ib;
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__2, i__3, ib,
da_ref(i, i ), ldda, t_ref(i), nb,
da_ref(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue);
}
/* Apply H to rows i:m of current block on the CPU */
lapackf77_cungqr(&i__2, &ib, &ib,
a_ref(i, i), &lda,
&tau[i], work, &lwork, &iinfo);
magma_csetmatrix_async( i__2, ib,
a_ref(i,i), lda,
da_ref(i,i), ldda, queue, &event );
/* Set rows 1:i-1 of current block to zero */
i__2 = i + ib;
magmablas_claset(MagmaFull, i, i__2 - i, c_zero, c_zero, da_ref(0,i), ldda, queue);
}
}
magma_cgetmatrix(m, n, da_ref(0, 0), ldda, a_ref(0, 0), lda, queue);
//cudaStreamDestroy(stream);
magma_free( da );
magma_free_cpu(work);
return *info;
} /* magma_cungqr */
extern "C" magma_int_t
magma_cgesv_rbt(
magma_bool_t ref, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *B, magma_int_t ldb,
magma_int_t *info)
{
/* Function Body */
*info = 0;
if ( ! (ref == MagmaTrue) &&
! (ref == MagmaFalse) ) {
*info = -1;
}
else if (n < 0) {
*info = -2;
} else if (nrhs < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (nrhs == 0 || n == 0)
return *info;
magma_int_t nn = n + ((4-(n % 4))%4);
magmaFloatComplex *dA, *hu, *hv, *db, *dAo, *dBo, *dwork;
magma_int_t n2;
magma_int_t iter;
n2 = nn*nn;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, n2 )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc( &db, nn*nrhs )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (ref == MagmaTrue) {
if (MAGMA_SUCCESS != magma_cmalloc( &dAo, n2 )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc( &dwork, nn*nrhs )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc( &dBo, nn*nrhs )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &hu, 2*nn )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (MAGMA_SUCCESS != magma_cmalloc_cpu( &hv, 2*nn )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmablas_claset(MagmaFull, nn, nn, MAGMA_C_ZERO, MAGMA_C_ONE, dA, nn);
/* Send matrix on the GPU*/
magma_csetmatrix(n, n, A, lda, dA, nn);
/* Send b on the GPU*/
magma_csetmatrix(n, nrhs, B, ldb, db, nn);
*info = magma_cgerbt_gpu(MagmaTrue, nn, nrhs, dA, nn, db, nn, hu, hv, info);
if (*info != MAGMA_SUCCESS) {
return *info;
}
if (ref == MagmaTrue) {
magma_ccopymatrix(nn, nn, dA, nn, dAo, nn);
magma_ccopymatrix(nn, nrhs, db, nn, dBo, nn);
}
/* Solve the system U^TAV.y = U^T.b on the GPU*/
magma_cgesv_nopiv_gpu( nn, nrhs, dA, nn, db, nn, info);
/* Iterative refinement */
if (ref == MagmaTrue) {
magma_cgerfs_nopiv_gpu(MagmaNoTrans, nn, nrhs, dAo, nn, dBo, nn, db, nn, dwork, dA, &iter, info);
}
//printf("iter = %d\n", iter);
/* The solution of A.x = b is Vy computed on the GPU */
magmaFloatComplex *dv;
if (MAGMA_SUCCESS != magma_cmalloc( &dv, 2*nn )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_csetvector(2*nn, hv, 1, dv, 1);
for(int i = 0; i < nrhs; i++) {
magmablas_cprbt_mv(nn, dv, db+(i*nn));
}
magma_cgetmatrix(n, nrhs, db, nn, B, ldb);
magma_free_cpu( hu);
magma_free_cpu( hv);
magma_free( dA );
magma_free( dv );
magma_free( db );
if (ref == MagmaTrue) {
magma_free( dAo );
magma_free( dBo );
magma_free( dwork );
}
return *info;
}
magma_int_t
magma_cmgenerator(
magma_int_t n,
magma_int_t offdiags,
magma_index_t *diag_offset,
magmaFloatComplex *diag_vals,
magma_c_matrix *A,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_c_matrix B={Magma_CSR};
B.val = NULL;
B.col = NULL;
B.row = NULL;
B.rowidx = NULL;
B.blockinfo = NULL;
B.diag = NULL;
B.dval = NULL;
B.dcol = NULL;
B.drow = NULL;
B.drowidx = NULL;
B.ddiag = NULL;
B.list = NULL;
B.dlist = NULL;
B.num_rows = n;
B.num_cols = n;
B.fill_mode = MagmaFull;
B.memory_location = Magma_CPU;
B.storage_type = Magma_ELLPACKT;
B.max_nnz_row = (2*offdiags+1);
CHECK( magma_cmalloc_cpu( &B.val, B.max_nnz_row*n ));
CHECK( magma_index_malloc_cpu( &B.col, B.max_nnz_row*n ));
for( int i=0; i<n; i++ ) { // stride over rows
// stride over the number of nonzeros in each row
// left of diagonal
for( int j=0; j<offdiags; j++ ) {
B.val[ i*B.max_nnz_row + j ] = diag_vals[ offdiags - j ];
B.col[ i*B.max_nnz_row + j ] = -1 * diag_offset[ offdiags-j ] + i;
}
// elements on the diagonal
B.val[ i*B.max_nnz_row + offdiags ] = diag_vals[ 0 ];
B.col[ i*B.max_nnz_row + offdiags ] = i;
// right of diagonal
for( int j=0; j<offdiags; j++ ) {
B.val[ i*B.max_nnz_row + j + offdiags +1 ] = diag_vals[ j+1 ];
B.col[ i*B.max_nnz_row + j + offdiags +1 ] = diag_offset[ j+1 ] + i;
}
}
// set invalid entries to zero
for( int i=0; i<n; i++ ) { // stride over rows
for( int j=0; j<B.max_nnz_row; j++ ) { // nonzeros in every row
if ( (B.col[i*B.max_nnz_row + j] < 0) ||
(B.col[i*B.max_nnz_row + j] >= n) ) {
B.val[ i*B.max_nnz_row + j ] = MAGMA_C_MAKE( 0.0, 0.0 );
}
}
}
B.nnz = 0;
for( int i=0; i<n; i++ ) { // stride over rows
for( int j=0; j<B.max_nnz_row; j++ ) { // nonzeros in every row
if ( MAGMA_C_REAL( B.val[i*B.max_nnz_row + j]) != 0.0 )
B.nnz++;
}
}
B.true_nnz = B.nnz;
// converting it to CSR will remove the invalit entries completely
CHECK( magma_cmconvert( B, A, Magma_ELLPACKT, Magma_CSR, queue ));
cleanup:
if( info != 0 ){
magma_cmfree( &B, queue );
}
return info;
}
/**
Purpose
-------
CGETRF_NOPIV_GPU computes an LU factorization of a general M-by-N
matrix A without any pivoting.
The factorization has the form
A = P * L * U
where P is a permutation matrix, 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]
dA COMPLEX array on the GPU, dimension (LDDA,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]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 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_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_nopiv_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define dA(i,j) (dA + (i)*nb + (j)*nb*ldda)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, mindim;
magma_int_t i, rows, s, lddwork;
magmaFloatComplex *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < 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;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
magma_cgetrf_nopiv( m, n, work, m, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( i=0; i < s; i++ ) {
// download i-th panel
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dA(i,i), ldda, work, lddwork, stream[0] );
if ( i > 0 ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n - (i+1)*nb,
c_one, dA(i-1,i-1), ldda,
dA(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-i*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i, i-1), ldda, dA(i-1,i+1), ldda,
c_one, dA(i, i+1), ldda );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
magma_cgetrf_nopiv( rows, nb, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dA(i, i), ldda, stream[0] );
magma_queue_sync( stream[0] );
// do the small non-parallel computations
if ( s > (i+1) ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n-s*nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
magma_cgetmatrix( rows, nb0, dA(s,s), ldda, work, lddwork );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
magma_cgetrf_nopiv( rows, nb0, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dA(s,s), ldda );
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb0, n-s*nb-nb0,
c_one, dA(s,s), ldda,
dA(s,s)+nb0, ldda);
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_cgetrf_nopiv_gpu */
/**
Purpose
-------
CGEHRD reduces a COMPLEX general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H . This version
stores the triangular matrices used in the factorization so that they can
be applied directly (i.e., without being recomputed) later. As a result,
the application of Q is much faster.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ilo INTEGER
@param[in]
ihi INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to CGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, and the
elements below the first subdiagonal, with the array TAU,
represent the orthogonal matrix Q as a product of elementary
reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
@param[out]
work (workspace) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= max(1,N).
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT COMPLEX array on the GPU, dimension NB*N,
where NB is the optimal blocksize. It stores the NB*NB blocks
of the triangular T matrices used in the reduction.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
---------------
The matrix Q is represented as a product of (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
@verbatim
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
@endverbatim
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
This version stores the T matrices in dT, for later use in magma_cunghr.
@ingroup magma_cgeev_comp
********************************************************************/
extern "C" magma_int_t
magma_cgehrd(
magma_int_t n, magma_int_t ilo, magma_int_t ihi,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magmaFloatComplex_ptr dT,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magma_int_t nb = magma_get_cgehrd_nb(n);
magma_int_t ldda = ((n+31)/32)*32;
magma_int_t i, nh, iws;
magma_int_t iinfo;
magma_int_t lquery;
*info = 0;
iws = n*nb;
work[0] = MAGMA_C_MAKE( iws, 0 );
lquery = (lwork == -1);
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
// Adjust from 1-based indexing
ilo -= 1;
// Quick return if possible
nh = ihi - ilo;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
// If not enough workspace, use unblocked code
if ( lwork < iws ) {
nb = 1;
}
if (nb == 1 || nb > nh) {
// Use unblocked code below
i = ilo;
}
else {
// Use blocked code
// GPU workspace is:
// nb*ldda for dwork for clahru
// nb*ldda for dV
// n*ldda for dA
magmaFloatComplex *dwork;
if (MAGMA_SUCCESS != magma_cmalloc( &dwork, 2*nb*ldda + n*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaFloatComplex *dV = dwork + nb*ldda;
magmaFloatComplex *dA = dwork + nb*ldda*2;
magmaFloatComplex *dTi;
magmaFloatComplex *T;
magma_cmalloc_cpu( &T, nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// zero first block of V, which is lower triangular
magmablas_claset( MagmaFull, nb, nb, c_zero, c_zero, dV, ldda );
// Set elements 0:ILO-1 and IHI-1:N-2 of TAU to zero
for (i = 0; i < ilo; ++i)
tau[i] = c_zero;
for (i = max(0,ihi-1); i < n-1; ++i)
tau[i] = c_zero;
assert( nb % 4 == 0 );
for (i=0; i < nb*nb; i += 4)
T[i] = T[i+1] = T[i+2] = T[i+3] = c_zero;
magmablas_claset( MagmaFull, nb, n, c_zero, c_zero, dT, nb );
// Copy the matrix to the GPU
magma_csetmatrix( n, n-ilo, A(0,ilo), lda, dA, ldda );
for (i = ilo; i < ihi-1 - nb; i += nb) {
// Reduce columns i:i+nb-1 to Hessenberg form, returning the
// matrices V and T of the block reflector H = I - V*T*V'
// which performs the reduction, and also the matrix Y = A*V*T
// Get the current panel (no need for the 1st iteration)
magma_cgetmatrix( ihi-i, nb,
dA(i,i-ilo), ldda,
A(i,i), lda );
// add 1 to i for 1-based index
magma_clahr2( ihi, i+1, nb,
dA(0,i-ilo), ldda,
dV, ldda,
A(0,i), lda,
&tau[i], T, nb, work, n);
// Copy T from the CPU to dT on the GPU
dTi = dT + (i - ilo)*nb;
magma_csetmatrix( nb, nb, T, nb, dTi, nb );
magma_clahru( n, ihi, i, nb,
A(0,i), lda,
dA(0,i-ilo), ldda, // dA
dA(i,i-ilo), ldda, // dY, stored over current panel
dV, ldda,
dTi, dwork );
}
// Copy remainder to host
magma_cgetmatrix( n, n-i,
dA(0,i-ilo), ldda,
A(0,i), lda );
magma_free( dwork );
magma_free_cpu( T );
}
// Use unblocked code to reduce the rest of the matrix
// add 1 to i for 1-based index
i += 1;
lapackf77_cgehd2(&n, &i, &ihi, A, &lda, tau, work, &iinfo);
work[0] = MAGMA_C_MAKE( iws, 0 );
return *info;
} /* magma_cgehrd */
extern "C" magma_int_t
magma_cbicgstab_merge(
magma_c_sparse_matrix A, magma_c_vector b,
magma_c_vector *x, magma_c_solver_par *solver_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_BICGSTABMERGE;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
// some useful variables
magmaFloatComplex c_zero = MAGMA_C_ZERO, c_one = MAGMA_C_ONE;
magma_int_t dofs = A.num_rows;
// GPU stream
magma_queue_t stream[2];
magma_event_t event[1];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_create( &event[0] );
// workspace
magma_c_vector q, r,rr,p,v,s,t;
magmaFloatComplex *d1, *d2, *skp;
d1 = NULL;
d2 = NULL;
skp = NULL;
magma_int_t stat_dev = 0, stat_cpu = 0;
stat_dev += magma_cmalloc( &d1, dofs*(2) );
stat_dev += magma_cmalloc( &d2, dofs*(2) );
// array for the parameters
stat_dev += magma_cmalloc( &skp, 8 );
if( stat_dev != 0 ){
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
printf("error: memory allocation.\n");
return MAGMA_ERR_DEVICE_ALLOC;
}
// skp = [alpha|beta|omega|rho_old|rho|nom|tmp1|tmp2]
magma_c_vinit( &q, Magma_DEV, dofs*6, c_zero, queue );
// q = rr|r|p|v|s|t
rr.memory_location = Magma_DEV; rr.dval = NULL; rr.num_rows = rr.nnz = dofs; rr.num_cols = 1;
r.memory_location = Magma_DEV; r.dval = NULL; r.num_rows = r.nnz = dofs; r.num_cols = 1;
p.memory_location = Magma_DEV; p.dval = NULL; p.num_rows = p.nnz = dofs; p.num_cols = 1;
v.memory_location = Magma_DEV; v.dval = NULL; v.num_rows = v.nnz = dofs; v.num_cols = 1;
s.memory_location = Magma_DEV; s.dval = NULL; s.num_rows = s.nnz = dofs; s.num_cols = 1;
t.memory_location = Magma_DEV; t.dval = NULL; t.num_rows = t.nnz = dofs; t.num_cols = 1;
rr.dval = q(0);
r.dval = q(1);
p.dval = q(2);
v.dval = q(3);
s.dval = q(4);
t.dval = q(5);
// solver variables
magmaFloatComplex alpha, beta, omega, rho_old, rho_new, *skp_h;
float nom, nom0, betanom, r0, den;
// solver setup
magma_cscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_ccopy( dofs, b.dval, 1, q(0), 1 ); // rr = b
magma_ccopy( dofs, b.dval, 1, q(1), 1 ); // r = b
rho_new = magma_cdotc( dofs, r.dval, 1, r.dval, 1 ); // rho=<rr,r>
nom = MAGMA_C_REAL(magma_cdotc( dofs, r.dval, 1, r.dval, 1 ));
nom0 = betanom = sqrt(nom); // nom = || r ||
rho_old = omega = alpha = MAGMA_C_MAKE( 1.0, 0. );
beta = rho_new;
solver_par->init_res = nom0;
// array on host for the parameters
stat_cpu = magma_cmalloc_cpu( &skp_h, 8 );
if( stat_cpu != 0 ){
magma_free( d1 );
magma_free( d2 );
magma_free( skp );
magma_free_cpu( skp_h );
printf("error: memory allocation.\n");
return MAGMA_ERR_HOST_ALLOC;
}
skp_h[0]=alpha;
skp_h[1]=beta;
skp_h[2]=omega;
skp_h[3]=rho_old;
skp_h[4]=rho_new;
skp_h[5]=MAGMA_C_MAKE(nom, 0.0);
magma_csetvector( 8, skp_h, 1, skp, 1 );
magma_c_spmv( c_one, A, r, c_zero, v, queue ); // z = A r
den = MAGMA_C_REAL( magma_cdotc(dofs, v.dval, 1, r.dval, 1) );// den = z dot r
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
// check positive definite
if (den <= 0.0) {
printf("Operator A is not postive definite. (Ar,r) = %f\n", den);
magmablasSetKernelStream( orig_queue );
return MAGMA_NONSPD;
solver_par->info = MAGMA_NONSPD;;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
magmablasSetKernelStream(stream[0]);
// computes p=r+beta*(p-omega*v)
magma_cbicgmerge1( dofs, skp, v.dval, r.dval, p.dval, queue );
magma_c_spmv( c_one, A, p, c_zero, v, queue ); // v = Ap
magma_cmdotc( dofs, 1, q.dval, v.dval, d1, d2, skp, queue );
magma_cbicgmerge4( 1, skp, queue );
magma_cbicgmerge2( dofs, skp, r.dval, v.dval, s.dval, queue ); // s=r-alpha*v
magma_c_spmv( c_one, A, s, c_zero, t, queue ); // t=As
magma_cmdotc( dofs, 2, q.dval+4*dofs, t.dval, d1, d2, skp+6, queue );
magma_cbicgmerge4( 2, skp, queue );
magma_cbicgmerge_xrbeta( dofs, d1, d2, q.dval, r.dval, p.dval,
s.dval, t.dval, x->dval, skp, queue );
// check stopping criterion (asynchronous copy)
magma_cgetvector_async( 1 , skp+5, 1,
skp_h+5, 1, stream[1] );
betanom = sqrt(MAGMA_C_REAL(skp_h[5]));
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) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < r0 ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_cresidual( A, b, *x, &residual, queue );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
magma_c_vfree(&q, queue ); // frees all vectors
magma_free(d1);
magma_free(d2);
magma_free( skp );
magma_free_cpu( skp_h );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* cbicgstab_merge */
/**
Purpose
-------
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
A COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
lda INTEGER
The first dimension of the array A. LDA >= max(1,M).
@param[in]
tau COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
@param[in]
T COMPLEX array, dimension (NB, min(M,N)).
T contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu (except stored on the CPU, not the GPU).
@param[in]
nb INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in T.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cungqr_m(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *T, magma_int_t nb,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(d,i,j) (dA[d] + (i) + (j)*ldda)
#define dT(d,i,j) (dT[d] + (i) + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldwork;
magma_int_t d, i, ib, j, jb, ki, kk;
magmaFloatComplex *work=NULL;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
magma_int_t di, dn;
magma_int_t dpanel;
magma_int_t ngpu = magma_num_gpus();
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// Allocate memory on GPUs for A and workspaces
magma_int_t ldda = magma_roundup( m, 32 );
magma_int_t lddwork = magma_roundup( n, 32 );
magma_int_t min_lblocks = (n / nb) / ngpu; // min. blocks per gpu
magma_int_t last_dev = (n / nb) % ngpu; // device with last block
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magmaFloatComplex *dA[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dT[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dV[ MagmaMaxGPUs ] = { NULL };
magmaFloatComplex *dW[ MagmaMaxGPUs ] = { NULL };
magma_queue_t queues[ MagmaMaxGPUs ] = { NULL };
for( d = 0; d < ngpu; ++d ) {
// example with n = 75, nb = 10, ngpu = 3
// min_lblocks = 2
// last_dev = 1
// gpu 0: 2 blocks, cols: 0- 9, 30-39, 60-69
// gpu 1: 1+ blocks, cols: 10-19, 40-49, 70-74 (partial)
// gpu 2: 1 block, cols: 20-29, 50-59
magma_setdevice( d );
nlocal[d] = min_lblocks*nb;
if ( d < last_dev ) {
nlocal[d] += nb;
}
else if ( d == last_dev ) {
nlocal[d] += (n % nb);
}
ldwork = nlocal[d]*ldda // dA
+ nb*m // dT
+ nb*ldda // dV
+ nb*lddwork; // dW
if ( MAGMA_SUCCESS != magma_cmalloc( &dA[d], ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto cleanup;
}
dT[d] = dA[d] + nlocal[d]*ldda;
dV[d] = dT[d] + nb*m;
dW[d] = dV[d] + nb*ldda;
magma_queue_create( d, &queues[d] );
}
trace_init( 1, ngpu, 1, queues );
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for larfb work
// m*nb for V
// nb*nb for T
lwork = (n + m + nb) * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
goto cleanup;
}
magmaFloatComplex *work_T, *work_V;
work_T = work + n*nb;
work_V = work + n*nb + nb*nb;
// Use unblocked code for the last or only block.
if (kk < n) {
trace_cpu_start( 0, "ungqr", "ungqr last block" );
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
// cungqr requires less workspace (n*nb), but is slow if k < cungqr's block size.
// replacing it with the 4 routines below is much faster (e.g., 60x).
//magma_int_t iinfo;
//lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
// A(kk, kk), &lda,
// &tau[kk], work, &lwork, &iinfo );
lapackf77_clacpy( MagmaFullStr, &m_kk, &k_kk, A(kk,kk), &lda, work_V, &m_kk);
lapackf77_claset( MagmaFullStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
work_V, &m_kk, &tau[kk], work_T, &k_kk);
lapackf77_clarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
work_V, &m_kk, work_T, &k_kk, A(kk, kk), &lda, work, &n_kk );
if (kk > 0) {
for( j=kk; j < n; j += nb ) {
jb = min( n-j, nb );
d = (j / nb) % ngpu;
di = ((j / nb) / ngpu) * nb;
magma_setdevice( d );
magma_csetmatrix( m_kk, jb,
A(kk, j), lda,
dA(d, kk, di), ldda, queues[d] );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaFull, kk, jb, c_zero, c_zero, dA(d, 0, di), ldda, queues[d] );
}
}
trace_cpu_end( 0 );
}
if (kk > 0) {
// Use blocked code
// send T to all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set T" );
magma_csetmatrix_async( nb, min(m,n), T, nb, dT[d], nb, queues[d] );
trace_gpu_end( d, 0 );
}
// queue: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
for( i = ki; i >= 0; i -= nb ) {
ib = min(nb, k - i);
mi = m - i;
dpanel = (i / nb) % ngpu;
di = ((i / nb) / ngpu) * nb;
// Send current panel to dV on the GPUs
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set V" );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV[d], ldda, queues[d] );
trace_gpu_end( d, 0 );
}
// set panel to identity
magma_setdevice( dpanel );
trace_gpu_start( dpanel, 0, "laset", "laset" );
magmablas_claset( MagmaFull, i, ib, c_zero, c_zero, dA(dpanel, 0, di), ldda, queues[dpanel] );
magmablas_claset( MagmaFull, mi, ib, c_zero, c_one, dA(dpanel, i, di), ldda, queues[dpanel] );
trace_gpu_end( dpanel, 0 );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_indices_1D_bcyclic( nb, ngpu, d, i, n, &di, &dn );
trace_gpu_start( d, 0, "larfb", "larfb" );
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, dn-di, ib,
dV[d], ldda, dT(d,0,i), nb,
dA(d, i, di), ldda, dW[d], lddwork, queues[d] );
trace_gpu_end( d, 0 );
}
}
}
// copy result back to CPU
trace_cpu_start( 0, "get", "get A" );
magma_cgetmatrix_1D_col_bcyclic( m, n, dA, ldda, A, lda, ngpu, nb, queues );
trace_cpu_end( 0 );
}
#ifdef TRACING
char name[80];
snprintf( name, sizeof(name), "cungqr-n%d-ngpu%d.svg", m, ngpu );
trace_finalize( name, "trace.css" );
#endif
cleanup:
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_free( dA[d] );
magma_queue_destroy( queues[d] );
}
magma_free_cpu( work );
magma_setdevice( orig_dev );
return *info;
} /* magma_cungqr */
extern "C" magma_int_t
magma_cgetrf_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
magmaFloatComplex **d_lA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, 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
=========
NUM_GPUS
(input) INTEGER
The number of GPUS to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array on the GPU, dimension (LDDA,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.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
> 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.
===================================================================== */
#define inAT(id,i,j) (d_lAT[(id)] + (i)*nb*lddat + (j)*nb)
magma_int_t nb, n_local[MagmaMaxGPUs];
magma_int_t maxm, mindim;
magma_int_t i, j, d, lddat, lddwork;
magmaFloatComplex *d_lAT[MagmaMaxGPUs];
magmaFloatComplex *d_panel[MagmaMaxGPUs], *work;
magma_queue_t streaml[MagmaMaxGPUs][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, d_lA[0], ldda, work, m );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, d_lA[0], ldda );
magma_free_cpu(work);
} else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if( num_gpus > ceil((float)n/nb) ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) num_gpus );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = ((((((n+nb-1)/nb)/num_gpus)*nb)+31)/32)*32;
lddat = (n+nb-1)/nb; /* number of block columns */
lddat = (lddat+num_gpus-1)/num_gpus; /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
lddat = ((lddat+31)/32)*32; /* make it a multiple of 32 */
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
/* local-n and local-ld */
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
/* workspaces */
if (MAGMA_SUCCESS != magma_cmalloc( &d_panel[i], (3+num_gpus)*nb*maxm )) {
for( j=0; j<=i; j++ ) {
magma_setdevice(j);
}
for( j=0; j<i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* local-matrix storage */
if (MAGMA_SUCCESS != magma_cmalloc( &d_lAT[i], lddat*maxm )) {
for( j=0; j<=i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
}
for( j=0; j<i; j++ ) {
magma_setdevice(j);
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* create the streams */
magma_queue_create( &streaml[i][0] );
magma_queue_create( &streaml[i][1] );
magmablasSetKernelStream(streaml[i][1]);
magmablas_ctranspose2( d_lAT[i], lddat, d_lA[i], ldda, m, n_local[i] );
}
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
cudaStreamSynchronize(streaml[i][0]);
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
/* cpu workspace */
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, lddwork*nb*num_gpus )) {
for(i=0; i<num_gpus; i++ ) {
magma_setdevice(i);
magma_free( d_panel[i] );
magma_free( d_lAT[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* calling multi-gpu interface with allocated workspaces and streams */
//magma_cgetrf1_mgpu( num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
// (magma_queue_t **)streaml, info );
magma_cgetrf2_mgpu(num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
streaml, info);
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
/* save on output */
magmablas_ctranspose2( d_lA[d], ldda, d_lAT[d], lddat, n_local[d], m );
magma_device_sync();
magma_free( d_lAT[d] );
magma_free( d_panel[d] );
magma_queue_destroy( streaml[d][0] );
magma_queue_destroy( streaml[d][1] );
magmablasSetKernelStream(NULL);
} /* end of for d=1,..,num_gpus */
magma_setdevice(0);
magma_free_pinned( work );
}
return *info;
}
/**
Purpose
-------
CHEEVDX_GPU computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm makes very mild assumptions about
floating point arithmetic. It will work on machines with a guard
digit in add/subtract, or on those binary machines without guard
digits which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or
Cray-2. It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@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,out]
dA COMPLEX array on the GPU,
dimension (LDDA, N).
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = MagmaVec, then if INFO = 0, the first m columns
of A contains the required
orthonormal eigenvectors of the matrix A.
If JOBZ = MagmaNoVec, then on exit the lower triangle (if UPLO=MagmaLower)
or the upper triangle (if UPLO=MagmaUpper) of A, including the
diagonal, is destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl REAL
@param[in]
vu REAL
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w REAL array, dimension (N)
If INFO = 0, the required m eigenvalues in ascending order.
@param
wA (workspace) COMPLEX array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LWORK >= N + N*NB.
If JOBZ = MagmaVec and N > 1, LWORK >= max( N + N*NB, 2*N + N**2 ).
NB can be obtained through magma_get_chetrd_nb(N).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK, RWORK and
IWORK arrays, returns these values as the first entries of
the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
rwork (workspace) REAL array, dimension (LRWORK)
On exit, if INFO = 0, RWORK[0] returns the optimal LRWORK.
@param[in]
lrwork INTEGER
The dimension of the array RWORK.
If N <= 1, LRWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LRWORK >= N.
If JOBZ = MagmaVec and N > 1, LRWORK >= 1 + 5*N + 2*N**2.
\n
If LRWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
iwork (workspace) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK[0] returns the optimal LIWORK.
@param[in]
liwork INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = MagmaNoVec and N > 1, LIWORK >= 1.
If JOBZ = MagmaVec and N > 1, LIWORK >= 3 + 5*N.
\n
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK, RWORK
and IWORK arrays, returns these values as the first entries
of the WORK, RWORK and IWORK arrays, and no error message
related to LWORK or LRWORK or LIWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i and JOBZ = MagmaNoVec, then the algorithm failed
to converge; i off-diagonal elements of an intermediate
tridiagonal form did not converge to zero;
if INFO = i and JOBZ = MagmaVec, then the algorithm failed
to compute an eigenvalue while working on the submatrix
lying in rows and columns INFO/(N+1) through
mod(INFO,N+1).
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified description of INFO. Sven, 16 Feb 05.
@ingroup magma_cheev_driver
********************************************************************/
extern "C" magma_int_t
magma_cheevdx_gpu(magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo,
magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *m, float *w,
magmaFloatComplex *wA, magma_int_t ldwa,
magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t lrwork,
magma_int_t *iwork, magma_int_t liwork,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
magma_int_t ione = 1;
float d__1;
float eps;
magma_int_t inde;
float anrm;
magma_int_t imax;
float rmin, rmax;
float sigma;
magma_int_t iinfo, lwmin;
magma_int_t lower;
magma_int_t llrwk;
magma_int_t wantz;
magma_int_t indwk2, llwrk2;
magma_int_t iscale;
float safmin;
float bignum;
magma_int_t indtau;
magma_int_t indrwk, indwrk, liwmin;
magma_int_t lrwmin, llwork;
float smlnum;
magma_int_t lquery;
magma_int_t alleig, valeig, indeig;
float *dwork;
magmaFloatComplex *dC;
magma_int_t lddc = ldda;
wantz = (jobz == MagmaVec);
lower = (uplo == MagmaLower);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1 || lrwork == -1 || liwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (ldwa < max(1,n)) {
*info = -14;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_chetrd_nb( n );
if ( n <= 1 ) {
lwmin = 1;
lrwmin = 1;
liwmin = 1;
}
else if ( wantz ) {
lwmin = max( n + n*nb, 2*n + n*n );
lrwmin = 1 + 5*n + 2*n*n;
liwmin = 3 + 5*n;
}
else {
lwmin = n + n*nb;
lrwmin = n;
liwmin = 1;
}
// multiply by 1+eps (in Double!) to ensure length gets rounded up,
// if it cannot be exactly represented in floating point.
real_Double_t one_eps = 1. + lapackf77_slamch("Epsilon");
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.);
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
if ((lwork < lwmin) && !lquery) {
*info = -16;
} else if ((lrwork < lrwmin) && ! lquery) {
*info = -18;
} else if ((liwork < liwmin) && ! lquery) {
*info = -20;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
}
else if (lquery) {
return *info;
}
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaFloatComplex *A;
magma_cmalloc_cpu( &A, n*n );
magma_cgetmatrix(n, n, dA, ldda, A, n);
lapackf77_cheevd(jobz_, uplo_,
&n, A, &n,
w, work, &lwork,
rwork, &lrwork,
iwork, &liwork, info);
magma_csetmatrix( n, n, A, n, dA, ldda);
magma_free_cpu(A);
*m=n;
return *info;
}
magma_queue_t stream;
magma_queue_create( &stream );
// dC and dwork are never used together, so use one buffer for both;
// unfortunately they're different types (complex and float).
// (this works better in dsyevd_gpu where they're both float).
// n*lddc for chetrd2_gpu, *2 for complex
// n for clanhe
magma_int_t ldwork = n*lddc*2;
if ( wantz ) {
// need 3n^2/2 for cstedx
ldwork = max( ldwork, 3*n*(n/2 + 1) );
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dC = (magmaFloatComplex*) dwork;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_clanhe(MagmaMaxNorm, uplo, n, dA, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
magmablas_clascl(uplo, 0, 0, 1., sigma, n, n, dA, ldda, info);
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
// chetrd rwork: e (n)
// cstedx rwork: e (n) + llrwk (1 + 4*N + 2*N**2) ==> 1 + 5n + 2n^2
inde = 0;
indrwk = inde + n;
llrwk = lrwork - indrwk;
// chetrd work: tau (n) + llwork (n*nb) ==> n + n*nb
// cstedx work: tau (n) + z (n^2)
// cunmtr work: tau (n) + z (n^2) + llwrk2 (n or n*nb) ==> 2n + n^2, or n + n*nb + n^2
indtau = 0;
indwrk = indtau + n;
indwk2 = indwrk + n*n;
llwork = lwork - indwrk;
llwrk2 = lwork - indwk2;
magma_timer_t time=0;
timer_start( time );
#ifdef FAST_HEMV
magma_chetrd2_gpu(uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork,
dC, n*lddc, &iinfo);
#else
magma_chetrd_gpu (uplo, n, dA, ldda, w, &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
timer_stop( time );
timer_printf( "time chetrd_gpu = %6.2f\n", time );
/* For eigenvalues only, call SSTERF. For eigenvectors, first call
CSTEDC to generate the eigenvector matrix, WORK(INDWRK), of the
tridiagonal matrix, then call CUNMTR to multiply it to the Householder
transformations represented as Householder vectors in A. */
if (! wantz) {
lapackf77_ssterf(&n, w, &rwork[inde], info);
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
}
else {
timer_start( time );
magma_cstedx(range, n, vl, vu, il, iu, w, &rwork[inde],
&work[indwrk], n, &rwork[indrwk],
llrwk, iwork, liwork, dwork, info);
timer_stop( time );
timer_printf( "time cstedx = %6.2f\n", time );
timer_start( time );
magma_smove_eig(range, n, w, &il, &iu, vl, vu, m);
magma_csetmatrix( n, *m, &work[indwrk + n * (il-1) ], n, dC, lddc );
magma_cunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dC, lddc, wA, ldwa, &iinfo);
magma_ccopymatrix( n, *m, dC, lddc, dA, ldda );
timer_stop( time );
timer_printf( "time cunmtr_gpu + copy = %6.2f\n", time );
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = n;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, w, &ione);
}
work[0] = MAGMA_C_MAKE( lwmin * one_eps, 0.); // round up
rwork[0] = lrwmin * one_eps;
iwork[0] = liwmin;
magma_queue_destroy( stream );
magma_free( dwork );
return *info;
} /* magma_cheevdx_gpu */
/**
Purpose
-------
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]
dA COMPLEX array on the GPU, 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 CGEQRF.
@param[in]
ldda INTEGER
The leading dimension of the array A, LDDA >= M.
@param[in,out]
dB COMPLEX 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]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) COMPLEX 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 >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by magma_get_cgeqrf_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 HWORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cgels_driver
********************************************************************/
extern "C" magma_int_t
magma_cgels_gpu( magma_trans_t trans, magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dB, magma_int_t lddb,
magmaFloatComplex *hwork, magma_int_t lwork,
magma_int_t *info)
{
magmaFloatComplex *dT;
magmaFloatComplex *tau;
magma_int_t k;
magma_int_t nb = magma_get_cgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_C_MAKE( (float)lwkopt, 0. );
*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 (ldda < max(1,m))
*info = -6;
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] = MAGMA_C_ONE;
return *info;
}
/*
* Allocate temporary buffers
*/
int ldtwork = ( 2*k + ((n+31)/32)*32 )*nb;
if (nb < nrhs)
ldtwork = ( 2*k + ((n+31)/32)*32 )*nrhs;
if (MAGMA_SUCCESS != magma_cmalloc( &dT, ldtwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_cmalloc_cpu( &tau, k );
if ( tau == NULL ) {
magma_free( dT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgeqrf_gpu( m, n, dA, ldda, tau, dT, info );
if ( *info == 0 ) {
magma_cgeqrs_gpu( m, n, nrhs,
dA, ldda, tau, dT,
dB, lddb, hwork, lwork, info );
}
magma_free( dT );
magma_free_cpu(tau);
return *info;
}
magma_int_t
magma_cm_5stencil(
magma_int_t n,
magma_c_matrix *A,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t i,j,k;
magma_c_matrix hA={Magma_CSR};
// generate matrix of desired structure and size (2d 5-point stencil)
magma_int_t nn = n*n;
magma_int_t offdiags = 2;
magma_index_t *diag_offset=NULL;
magmaFloatComplex *diag_vals=NULL;
CHECK( magma_cmalloc_cpu( &diag_vals, offdiags+1 ));
CHECK( magma_index_malloc_cpu( &diag_offset, offdiags+1 ));
diag_offset[0] = 0;
diag_offset[1] = 1;
diag_offset[2] = n;
#define COMPLEX
#ifdef COMPLEX
// complex case
diag_vals[0] = MAGMA_C_MAKE( 4.0, 4.0 );
diag_vals[1] = MAGMA_C_MAKE( -1.0, -1.0 );
diag_vals[2] = MAGMA_C_MAKE( -1.0, -1.0 );
#else
// real case
diag_vals[0] = MAGMA_C_MAKE( 4.0, 0.0 );
diag_vals[1] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[2] = MAGMA_C_MAKE( -1.0, 0.0 );
#endif
CHECK( magma_cmgenerator( nn, offdiags, diag_offset, diag_vals, &hA, queue ));
// now set some entries to zero (boundary...)
for( i=0; i<n; i++ ) {
for( j=0; j<n; j++ ) {
magma_index_t row = i*n+j;
for( k=hA.row[row]; k<hA.row[row+1]; k++) {
if ((hA.col[k] == row-1 ) && (row+1)%n == 1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
if ((hA.col[k] == row+1 ) && (row)%n ==n-1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
}
}
}
CHECK( magma_cmconvert( hA, A, Magma_CSR, Magma_CSR, queue ));
magma_cmcsrcompressor( A, queue );
A->true_nnz = A->nnz;
cleanup:
magma_free_cpu( diag_vals );
magma_free_cpu( diag_offset );
magma_cmfree( &hA, queue );
return info;
}
extern "C" magma_int_t
magma_ccsrsplit(
magma_int_t bsize,
magma_c_matrix A,
magma_c_matrix *D,
magma_c_matrix *R,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t i, k, j, nnz_diag, nnz_offd;
D->val = NULL;
D->col = NULL;
D->row = NULL;
D->rowidx = NULL;
D->blockinfo = NULL;
D->diag = NULL;
D->dval = NULL;
D->dcol = NULL;
D->drow = NULL;
D->drowidx = NULL;
D->ddiag = NULL;
R->val = NULL;
R->col = NULL;
R->row = NULL;
R->rowidx = NULL;
R->blockinfo = NULL;
R->diag = NULL;
R->dval = NULL;
R->dcol = NULL;
R->drow = NULL;
R->drowidx = NULL;
R->ddiag = NULL;
if ( A.memory_location == Magma_CPU &&
( A.storage_type == Magma_CSR ||
A.storage_type == Magma_CSRCOO ) ) {
nnz_diag = nnz_offd = 0;
// Count the new number of nonzeroes in the two matrices
for( i=0; i<A.num_rows; i+=bsize ){
for( k=i; k<min(A.num_rows,i+bsize); k++ ){
int check = 0;
for( j=A.row[k]; j<A.row[k+1]; j++ ){
if ( A.col[j] < i )
nnz_offd++;
else if ( A.col[j] < i+bsize ){
if( A.col[j] == k ){
check = 1;
}
nnz_diag++;
}
else
nnz_offd++;
}
if( check == 0 ){
printf("error: matrix contains zero on diagonal at (%d,%d).\n", i, i);
info = -1;
goto cleanup;
}
}
}
// Allocate memory for the new matrices
D->storage_type = Magma_CSRD;
D->memory_location = A.memory_location;
D->num_rows = A.num_rows;
D->num_cols = A.num_cols;
D->nnz = nnz_diag;
R->storage_type = Magma_CSR;
R->memory_location = A.memory_location;
R->num_rows = A.num_rows;
R->num_cols = A.num_cols;
R->nnz = nnz_offd;
CHECK( magma_cmalloc_cpu( &D->val, nnz_diag ));
CHECK( magma_index_malloc_cpu( &D->row, A.num_rows+1 ));
CHECK( magma_index_malloc_cpu( &D->col, nnz_diag ));
CHECK( magma_cmalloc_cpu( &R->val, nnz_offd ));
CHECK( magma_index_malloc_cpu( &R->row, A.num_rows+1 ));
CHECK( magma_index_malloc_cpu( &R->col, nnz_offd ));
// Fill up the new sparse matrices
D->row[0] = 0;
R->row[0] = 0;
nnz_offd = nnz_diag = 0;
for( i=0; i<A.num_rows; i+=bsize) {
for( k=i; k<min(A.num_rows,i+bsize); k++ ) {
D->row[k+1] = D->row[k];
R->row[k+1] = R->row[k];
for( j=A.row[k]; j<A.row[k+1]; j++ ) {
if ( A.col[j] < i ) {
R->val[nnz_offd] = A.val[j];
R->col[nnz_offd] = A.col[j];
R->row[k+1]++;
nnz_offd++;
}
else if ( A.col[j] < i+bsize ) {
// larger than diagonal remain as before
if ( A.col[j]>k ) {
D->val[nnz_diag] = A.val[ j ];
D->col[nnz_diag] = A.col[ j ];
D->row[k+1]++;
}
// diagonal is written first
else if ( A.col[j]==k ) {
D->val[D->row[k]] = A.val[ j ];
D->col[D->row[k]] = A.col[ j ];
D->row[k+1]++;
}
// smaller than diagonal are shifted one to the right
// to have room for the diagonal
else {
D->val[nnz_diag+1] = A.val[ j ];
D->col[nnz_diag+1] = A.col[ j ];
D->row[k+1]++;
}
nnz_diag++;
}
else {
R->val[nnz_offd] = A.val[j];
R->col[nnz_offd] = A.col[j];
R->row[k+1]++;
nnz_offd++;
}
}
}
}
}
else {
magma_c_matrix Ah={Magma_CSR}, ACSR={Magma_CSR}, DCSR={Magma_CSR}, RCSR={Magma_CSR}, Dh={Magma_CSR}, Rh={Magma_CSR};
CHECK( magma_cmtransfer( A, &Ah, A.memory_location, Magma_CPU, queue ));
CHECK( magma_cmconvert( Ah, &ACSR, A.storage_type, Magma_CSR, queue ));
CHECK( magma_ccsrsplit( bsize, ACSR, &DCSR, &RCSR, queue ));
CHECK( magma_cmconvert( DCSR, &Dh, Magma_CSR, A.storage_type, queue ));
CHECK( magma_cmconvert( RCSR, &Rh, Magma_CSR, A.storage_type, queue ));
CHECK( magma_cmtransfer( Dh, D, Magma_CPU, A.memory_location, queue ));
CHECK( magma_cmtransfer( Rh, R, Magma_CPU, A.memory_location, queue ));
magma_cmfree( &Ah, queue );
magma_cmfree( &ACSR, queue );
magma_cmfree( &Dh, queue );
magma_cmfree( &DCSR, queue );
magma_cmfree( &Rh, queue );
magma_cmfree( &RCSR, queue );
}
cleanup:
if( info != 0 ){
magma_cmfree( D, queue );
magma_cmfree( R, queue );
}
return info;
}
magma_int_t
magma_cm_27stencil(
magma_int_t n,
magma_c_matrix *A,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t i,j,k;
magma_c_matrix hA={Magma_CSR};
// generate matrix of desired structure and size (3d 27-point stencil)
magma_int_t nn = n*n*n;
magma_int_t offdiags = 13;
magma_index_t *diag_offset=NULL;
magmaFloatComplex *diag_vals=NULL;
CHECK( magma_cmalloc_cpu( &diag_vals, offdiags+1 ));
CHECK( magma_index_malloc_cpu( &diag_offset, offdiags+1 ));
diag_offset[0] = 0;
diag_offset[1] = 1;
diag_offset[2] = n-1;
diag_offset[3] = n;
diag_offset[4] = n+1;
diag_offset[5] = n*n-n-1;
diag_offset[6] = n*n-n;
diag_offset[7] = n*n-n+1;
diag_offset[8] = n*n-1;
diag_offset[9] = n*n;
diag_offset[10] = n*n+1;
diag_offset[11] = n*n+n-1;
diag_offset[12] = n*n+n;
diag_offset[13] = n*n+n+1;
diag_vals[0] = MAGMA_C_MAKE( 26.0, 0.0 );
diag_vals[1] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[2] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[3] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[4] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[5] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[6] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[7] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[8] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[9] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[10] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[11] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[12] = MAGMA_C_MAKE( -1.0, 0.0 );
diag_vals[13] = MAGMA_C_MAKE( -1.0, 0.0 );
CHECK( magma_cmgenerator( nn, offdiags, diag_offset, diag_vals, &hA, queue ));
// now set some entries to zero (boundary...)
for( i=0; i < n*n; i++ ) {
for( j=0; j < n; j++ ) {
magma_index_t row = i*n+j;
for( k=hA.row[row]; k<hA.row[row+1]; k++) {
if ((hA.col[k] == row-1 ||
hA.col[k] == row-n-1 ||
hA.col[k] == row+n-1 ||
hA.col[k] == row-n*n+n-1 ||
hA.col[k] == row+n*n-n-1 ||
hA.col[k] == row-n*n-1 ||
hA.col[k] == row+n*n-1 ||
hA.col[k] == row-n*n-n-1 ||
hA.col[k] == row+n*n+n-1 ) && (row+1)%n == 1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
if ((hA.col[k] == row+1 ||
hA.col[k] == row-n+1 ||
hA.col[k] == row+n+1 ||
hA.col[k] == row-n*n+n+1 ||
hA.col[k] == row+n*n-n+1 ||
hA.col[k] == row-n*n+1 ||
hA.col[k] == row+n*n+1 ||
hA.col[k] == row-n*n-n+1 ||
hA.col[k] == row+n*n+n+1 ) && (row)%n ==n-1 )
hA.val[k] = MAGMA_C_MAKE( 0.0, 0.0 );
}
}
}
hA.true_nnz = hA.nnz;
CHECK( magma_cmconvert( hA, A, Magma_CSR, Magma_CSR, queue ));
cleanup:
magma_free_cpu( diag_vals );
magma_free_cpu( diag_offset );
magma_cmfree( &hA, queue );
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 CGETRF_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 array on the GPU, dimension (LDA,N)
The factors L and U from the factorization A = P*L*U as computed
by CGETRF_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 CGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[in,out]
dB COMPLEX 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_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrs_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex_ptr dA, magma_int_t ldda, magma_int_t *ipiv,
magmaFloatComplex_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *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_cmalloc_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_cgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_claswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_csetmatrix( n, nrhs, work, n, dB, lddb );
if ( nrhs == 1) {
magma_ctrsv(MagmaLower, MagmaNoTrans, MagmaUnit, n, dA, ldda, dB, 1 );
magma_ctrsv(MagmaUpper, MagmaNoTrans, MagmaNonUnit, n, dA, ldda, dB, 1 );
} else {
magma_ctrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ctrsm(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_ctrsv(MagmaUpper, trans, MagmaNonUnit, n, dA, ldda, dB, 1 );
magma_ctrsv(MagmaLower, trans, MagmaUnit, n, dA, ldda, dB, 1 );
} else {
magma_ctrsm(MagmaLeft, MagmaUpper, trans, MagmaNonUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
magma_ctrsm(MagmaLeft, MagmaLower, trans, MagmaUnit, n, nrhs, c_one, dA, ldda, dB, lddb );
}
magma_cgetmatrix( n, nrhs, dB, lddb, work, n );
lapackf77_claswp(&nrhs, work, &n, &i1, &i2, ipiv, &inc);
magma_csetmatrix( n, nrhs, work, n, dB, lddb );
}
magma_free_cpu(work);
return *info;
}
static void magma_ctile_bulge_applyQ(magma_int_t core_id, char side, magma_int_t n_loc, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
magmaFloatComplex *E, magma_int_t lde, magmaFloatComplex *V, magma_int_t ldv,
magmaFloatComplex *TAU, magmaFloatComplex *T, magma_int_t ldt)//, magma_int_t* info)
{
//%===========================
//% local variables
//%===========================
magma_int_t firstcolj;
magma_int_t bg, rownbm;
magma_int_t st,ed,fst,vlen,vnb,colj;
magma_int_t vpos,tpos;
magma_int_t cur_blksiz,avai_blksiz, ncolinvolvd;
magma_int_t nbgr, colst, coled;
if(n<=0)
return ;
if(n_loc<=0)
return ;
//info = 0;
magma_int_t INFO=0;
magma_int_t nbGblk = magma_ceildiv(n-1, Vblksiz);
/*
* version v1: for each chunck it apply all the V's then move to
* the other chunck. the locality here inside each
* chunck meaning that thread t apply V_k then move
* to V_k+1 which overlap with V_k meaning that the
* E_k+1 overlap with E_k. so here is the
* locality however thread t had to read V_k+1 and
* T_k+1 at each apply. note that all thread if they
* run at same speed they might reading the same V_k
* and T_k at the same time.
* */
magma_int_t nb_loc = 128; //$$$$$$$$
magma_int_t lwork = 2*nb_loc*max(Vblksiz,64);
magmaFloatComplex *work, *work2;
magma_cmalloc_cpu(&work, lwork);
magma_cmalloc_cpu(&work2, lwork);
magma_int_t nbchunk = magma_ceildiv(n_loc, nb_loc);
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* each q_i consist of applying V to a block of row E(row_i,:) and applies are overlapped meaning
* that q_i+1 overlap a portion of the E(row_i, :).
* IN parallel E is splitten in vertical block over the threads */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* each q_i consist of applying V to a block of col E(:, col_i,:) and the applies are overlapped meaning
* that q_i+1 overlap a portion of the E(:, col_i).
* IN parallel E is splitten in horizontal block over the threads */
#ifdef ENABLE_DEBUG
if((core_id==0)||(core_id==1))
printf(" APPLY Q2_cpu cbulge_back N %d N_loc %d nbchunk %d NB %d Vblksiz %d SIDE %c \n", n, n_loc, nbchunk, nb, Vblksiz, side);
#endif
for (magma_int_t i = 0; i<nbchunk; i++)
{
magma_int_t ib_loc = min(nb_loc, (n_loc - i*nb_loc));
if(side=='L') {
for (bg = nbGblk; bg>0; bg--)
{
firstcolj = (bg-1)*Vblksiz + 1;
rownbm = magma_ceildiv((n-(firstcolj+1)),nb);
if(bg==nbGblk) rownbm = magma_ceildiv((n-(firstcolj)),nb); // last blk has size=1 used for complex to handle A(N,N-1)
for (magma_int_t j = rownbm; j>0; j--)
{
vlen = 0;
vnb = 0;
colj = (bg-1)*Vblksiz; // for k=0;I compute the fst and then can remove it from the loop
fst = (rownbm -j)*nb+colj +1;
for (magma_int_t k=0; k<Vblksiz; k++)
{
colj = (bg-1)*Vblksiz + k;
st = (rownbm -j)*nb+colj +1;
ed = min(st+nb-1,n-1);
if(st>ed)
break;
if((st==ed)&&(colj!=n-2))
break;
vlen=ed-fst+1;
vnb=k+1;
}
colst = (bg-1)*Vblksiz;
magma_bulge_findVTpos(n, nb, Vblksiz, colst, fst, ldv, ldt, &vpos, &tpos);
if((vlen>0)&&(vnb>0)) {
lapackf77_clarfb( "L", "N", "F", "C", &vlen, &ib_loc, &vnb, V(vpos), &ldv, T(tpos), &ldt, E(fst,i*nb_loc), &lde, work, &ib_loc);
}
if(INFO!=0)
printf("ERROR CUNMQR INFO %d \n", (int) INFO);
}
}
} else if (side=='R') {
rownbm = magma_ceildiv((n-1),nb);
for (magma_int_t k = 1; k<=rownbm; k++)
{
ncolinvolvd = min(n-1, k*nb);
avai_blksiz=min(Vblksiz,ncolinvolvd);
nbgr = magma_ceildiv(ncolinvolvd,avai_blksiz);
for (magma_int_t j = 1; j<=nbgr; j++)
{
vlen = 0;
vnb = 0;
cur_blksiz = min(ncolinvolvd-(j-1)*avai_blksiz, avai_blksiz);
colst = (j-1)*avai_blksiz;
coled = colst + cur_blksiz -1;
fst = (rownbm -k)*nb+colst +1;
for (colj=colst; colj<=coled; colj++)
{
st = (rownbm -k)*nb+colj +1;
ed = min(st+nb-1,n-1);
if(st>ed)
break;
if((st==ed)&&(colj!=n-2))
break;
vlen=ed-fst+1;
vnb=vnb+1;
}
magma_bulge_findVTpos(n, nb, Vblksiz, colst, fst, ldv, ldt, &vpos, &tpos);
if((vlen>0)&&(vnb>0)) {
lapackf77_clarfb( "R", "N", "F", "C", &ib_loc, &vlen, &vnb, V(vpos), &ldv, T(tpos), &ldt, E(i*nb_loc,fst), &lde, work, &ib_loc);
}
}
}
} else {
printf("ERROR SIDE %d \n",side);
}
} // END loop over the chunks
magma_free_cpu(work);
magma_free_cpu(work2);
}
/**
Purpose
-------
CLATRD2 reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, CLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, CLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by CHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_cheev_aux
********************************************************************/
extern "C" magma_int_t
magma_clatrd2(magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
float *e, magmaFloatComplex *tau,
magmaFloatComplex *W, magma_int_t ldw,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dW, magma_int_t lddw,
magmaFloatComplex *dwork, magma_int_t ldwork)
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
#define dW(i, j) (dW + (j)*lddw + (i))
magma_int_t i;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex value = MAGMA_C_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaFloatComplex alpha;
magmaFloatComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_cmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside clatrd
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_cgemv("No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_clacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_cgemv("No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, A(i, i+1), &ldw);
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_clarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_C_REAL( alpha );
*A(i-1,i) = MAGMA_C_MAKE( 1, 0 );
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_csetvector( i, A(0, i), 1, dA(0, i), 1 );
//#if (GPUSHMEM < 200)
//magma_chemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
// dA(0, i), ione, c_zero, dW(0, iw), ione);
//#else
magmablas_chemv_work(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_cgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_cgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( stream );
if (i < n-1) {
blasf77_cgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_cgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_cgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_cscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_cdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_caxpy(&i, &alpha, A(0, i), &ione,
W(0, iw), &ione);
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, W(i, 0), &ldw);
#endif
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, W(i, 0), &ldw);
lapackf77_clacgv(&i, A(i, 0), &lda);
#endif
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, A(i, 0), &lda);
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_clarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_C_REAL( alpha );
*A(i+1,i) = MAGMA_C_MAKE( 1, 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_csetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
//#if (GPUSHMEM < 200)
//magma_chemv(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
// dW(i+1, i), ione);
//#else
magmablas_chemv_work(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_cgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_cgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_cgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
// 3. Here is where we need it
magma_queue_sync( stream );
if (i != 0)
blasf77_caxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_cscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_cdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_caxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* magma_clatrd */
/**
Purpose
-------
CHEEVX computes selected eigenvalues and, optionally, eigenvectors
of a complex Hermitian matrix A. Eigenvalues and eigenvectors can
be selected by specifying either a range of values or a range of
indices for the desired eigenvalues.
Arguments
---------
@param[in]
jobz magma_vec_t
- = MagmaNoVec: Compute eigenvalues only;
- = MagmaVec: Compute eigenvalues and eigenvectors.
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
@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,out]
dA COMPLEX array, dimension (LDDA, N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the
leading N-by-N upper triangular part of A contains the
upper triangular part of the matrix A. If UPLO = MagmaLower,
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, the lower triangle (if UPLO=MagmaLower) or the upper
triangle (if UPLO=MagmaUpper) of A, including the diagonal, is
destroyed.
@param[in]
ldda INTEGER
The leading dimension of the array DA. LDDA >= max(1,N).
@param[in]
vl REAL
@param[in]
vu REAL
If RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
If RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[in]
abstol REAL
The absolute error tolerance for the eigenvalues.
An approximate eigenvalue is accepted as converged
when it is determined to lie in an interval [a,b]
of width less than or equal to
ABSTOL + EPS * max( |a|,|b| ),
\n
where EPS is the machine precision. If ABSTOL is less than
or equal to zero, then EPS*|T| will be used in its place,
where |T| is the 1-norm of the tridiagonal matrix obtained
by reducing A to tridiagonal form.
\n
Eigenvalues will be computed most accurately when ABSTOL is
set to twice the underflow threshold 2*SLAMCH('S'), not zero.
If this routine returns with INFO > 0, indicating that some
eigenvectors did not converge, try setting ABSTOL to
2*SLAMCH('S').
\n
See "Computing Small Singular Values of Bidiagonal Matrices
with Guaranteed High Relative Accuracy," by Demmel and
Kahan, LAPACK Working Note #3.
@param[out]
m INTEGER
The total number of eigenvalues found. 0 <= M <= N.
If RANGE = MagmaRangeAll, M = N, and if RANGE = MagmaRangeI, M = IU-IL+1.
@param[out]
w REAL array, dimension (N)
On normal exit, the first M elements contain the selected
eigenvalues in ascending order.
@param[out]
dZ COMPLEX array, dimension (LDDZ, max(1,M))
If JOBZ = MagmaVec, then if INFO = 0, the first M columns of Z
contain the orthonormal eigenvectors of the matrix A
corresponding to the selected eigenvalues, with the i-th
column of Z holding the eigenvector associated with W(i).
If an eigenvector fails to converge, then that column of Z
contains the latest approximation to the eigenvector, and the
index of the eigenvector is returned in IFAIL.
If JOBZ = MagmaNoVec, then Z is not referenced.
Note: the user must ensure that at least max(1,M) columns are
supplied in the array Z; if RANGE = MagmaRangeV, the exact value of M
is not known in advance and an upper bound must be used.
********* (workspace) If FAST_HEMV is defined DZ should be (LDDZ, max(1,N)) in both cases.
@param[in]
lddz INTEGER
The leading dimension of the array DZ. LDDZ >= 1, and if
JOBZ = MagmaVec, LDDZ >= max(1,N).
@param
wA (workspace) COMPLEX array, dimension (LDWA, N)
@param[in]
ldwa INTEGER
The leading dimension of the array wA. LDWA >= max(1,N).
@param
wZ (workspace) COMPLEX array, dimension (LDWZ, max(1,M))
@param[in]
ldwz INTEGER
The leading dimension of the array wZ. LDWZ >= 1, and if
JOBZ = MagmaVec, LDWZ >= max(1,N).
@param[out]
work (workspace) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= (NB+1)*N,
where NB is the max of the blocksize for CHETRD.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param
rwork (workspace) REAL array, dimension (7*N)
@param
iwork (workspace) INTEGER array, dimension (5*N)
@param[out]
ifail INTEGER array, dimension (N)
If JOBZ = MagmaVec, then if INFO = 0, the first M elements of
IFAIL are zero. If INFO > 0, then IFAIL contains the
indices of the eigenvectors that failed to converge.
If JOBZ = MagmaNoVec, then IFAIL is not referenced.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, then i eigenvectors failed to converge.
Their indices are stored in array IFAIL.
@ingroup magma_cheev_driver
********************************************************************/
extern "C" magma_int_t
magma_cheevx_gpu(
magma_vec_t jobz, magma_range_t range, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
float vl, float vu,
magma_int_t il, magma_int_t iu, float abstol, magma_int_t *m,
float *w,
magmaFloatComplex_ptr dZ, magma_int_t lddz,
magmaFloatComplex *wA, magma_int_t ldwa,
magmaFloatComplex *wZ, magma_int_t ldwz,
magmaFloatComplex *work, magma_int_t lwork,
float *rwork, magma_int_t *iwork, magma_int_t *ifail,
magma_int_t *info)
{
const char* uplo_ = lapack_uplo_const( uplo );
const char* jobz_ = lapack_vec_const( jobz );
const char* range_ = lapack_range_const( range );
magma_int_t ione = 1;
const char* order_;
magma_int_t indd, inde;
magma_int_t imax;
magma_int_t lopt, itmp1, indee;
magma_int_t lower, wantz;
magma_int_t i, j, jj, i__1;
magma_int_t alleig, valeig, indeig;
magma_int_t iscale, indibl;
magma_int_t indiwk, indisp, indtau;
magma_int_t indrwk, indwrk;
magma_int_t llwork, nsplit;
magma_int_t lquery;
magma_int_t iinfo;
float safmin;
float bignum;
float smlnum;
float eps, tmp1;
float anrm;
float sigma, d__1;
float rmin, rmax;
magmaFloat_ptr dwork;
/* Function Body */
lower = (uplo == MagmaLower);
wantz = (jobz == MagmaVec);
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
lquery = (lwork == -1);
*info = 0;
if (! (wantz || (jobz == MagmaNoVec))) {
*info = -1;
} else if (! (alleig || valeig || indeig)) {
*info = -2;
} else if (! (lower || (uplo == MagmaUpper))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (ldda < max(1,n)) {
*info = -6;
} else if (lddz < 1 || (wantz && lddz < n)) {
*info = -15;
} else if (ldwa < max(1,n)) {
*info = -17;
} else if (ldwz < 1 || (wantz && ldwz < n)) {
*info = -19;
} else {
if (valeig) {
if (n > 0 && vu <= vl) {
*info = -8;
}
} else if (indeig) {
if (il < 1 || il > max(1,n)) {
*info = -9;
} else if (iu < min(n,il) || iu > n) {
*info = -10;
}
}
}
magma_int_t nb = magma_get_chetrd_nb(n);
lopt = n * (nb + 1);
work[0] = MAGMA_C_MAKE( lopt, 0 );
if (lwork < lopt && ! lquery) {
*info = -21;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info));
return *info;
} else if (lquery) {
return *info;
}
*m = 0;
/* Check if matrix is very small then just call LAPACK on CPU, no need for GPU */
if (n <= 128) {
#ifdef ENABLE_DEBUG
printf("--------------------------------------------------------------\n");
printf(" warning matrix too small N=%d NB=%d, calling lapack on CPU \n", (int) n, (int) nb);
printf("--------------------------------------------------------------\n");
#endif
magmaFloatComplex *a;
magma_cmalloc_cpu( &a, n*n );
magma_cgetmatrix(n, n, dA, ldda, a, n);
lapackf77_cheevx(jobz_, range_, uplo_,
&n, a, &n, &vl, &vu, &il, &iu, &abstol, m,
w, wZ, &ldwz, work, &lwork,
rwork, iwork, ifail, info);
magma_csetmatrix( n, n, a, n, dA, ldda);
magma_csetmatrix( n, *m, wZ, ldwz, dZ, lddz);
magma_free_cpu(a);
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork, n )) {
fprintf (stderr, "!!!! device memory allocation error (magma_cheevx_gpu)\n");
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
--w;
--work;
--rwork;
--iwork;
--ifail;
/* Get machine constants. */
safmin = lapackf77_slamch("Safe minimum");
eps = lapackf77_slamch("Precision");
smlnum = safmin / eps;
bignum = 1. / smlnum;
rmin = magma_ssqrt(smlnum);
rmax = magma_ssqrt(bignum);
/* Scale matrix to allowable range, if necessary. */
anrm = magmablas_clanhe(MagmaMaxNorm, uplo, n, dA, ldda, dwork);
iscale = 0;
sigma = 1;
if (anrm > 0. && anrm < rmin) {
iscale = 1;
sigma = rmin / anrm;
} else if (anrm > rmax) {
iscale = 1;
sigma = rmax / anrm;
}
if (iscale == 1) {
d__1 = 1.;
magmablas_clascl(uplo, 0, 0, 1., sigma, n, n, dA,
ldda, info);
if (abstol > 0.) {
abstol *= sigma;
}
if (valeig) {
vl *= sigma;
vu *= sigma;
}
}
/* Call CHETRD to reduce Hermitian matrix to tridiagonal form. */
indd = 1;
inde = indd + n;
indrwk = inde + n;
indtau = 1;
indwrk = indtau + n;
llwork = lwork - indwrk + 1;
#ifdef FAST_HEMV
magma_chetrd2_gpu(uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, dZ, lddz*n, &iinfo);
#else
magma_chetrd_gpu (uplo, n, dA, ldda, &rwork[indd], &rwork[inde],
&work[indtau], wA, ldwa, &work[indwrk], llwork, &iinfo);
#endif
lopt = n + (magma_int_t)MAGMA_C_REAL(work[indwrk]);
/* If all eigenvalues are desired and ABSTOL is less than or equal to
zero, then call SSTERF or CUNGTR and CSTEQR. If this fails for
some eigenvalue, then try SSTEBZ. */
if ((alleig || (indeig && il == 1 && iu == n)) && abstol <= 0.) {
blasf77_scopy(&n, &rwork[indd], &ione, &w[1], &ione);
indee = indrwk + 2*n;
if (! wantz) {
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_ssterf(&n, &w[1], &rwork[indee], info);
}
else {
lapackf77_clacpy("A", &n, &n, wA, &ldwa, wZ, &ldwz);
lapackf77_cungtr(uplo_, &n, wZ, &ldwz, &work[indtau], &work[indwrk], &llwork, &iinfo);
i__1 = n - 1;
blasf77_scopy(&i__1, &rwork[inde], &ione, &rwork[indee], &ione);
lapackf77_csteqr(jobz_, &n, &w[1], &rwork[indee], wZ, &ldwz, &rwork[indrwk], info);
if (*info == 0) {
for (i = 1; i <= n; ++i) {
ifail[i] = 0;
}
magma_csetmatrix( n, n, wZ, ldwz, dZ, lddz );
}
}
if (*info == 0) {
*m = n;
}
}
/* Otherwise, call SSTEBZ and, if eigenvectors are desired, CSTEIN. */
if (*m == 0) {
*info = 0;
if (wantz) {
order_ = "B";
} else {
order_ = "E";
}
indibl = 1;
indisp = indibl + n;
indiwk = indisp + n;
lapackf77_sstebz(range_, order_, &n, &vl, &vu, &il, &iu, &abstol, &rwork[indd], &rwork[inde], m,
&nsplit, &w[1], &iwork[indibl], &iwork[indisp], &rwork[indrwk], &iwork[indiwk], info);
if (wantz) {
lapackf77_cstein(&n, &rwork[indd], &rwork[inde], m, &w[1], &iwork[indibl], &iwork[indisp],
wZ, &ldwz, &rwork[indrwk], &iwork[indiwk], &ifail[1], info);
magma_csetmatrix( n, *m, wZ, ldwz, dZ, lddz );
/* Apply unitary matrix used in reduction to tridiagonal
form to eigenvectors returned by CSTEIN. */
magma_cunmtr_gpu(MagmaLeft, uplo, MagmaNoTrans, n, *m, dA, ldda, &work[indtau],
dZ, lddz, wA, ldwa, &iinfo);
}
}
/* If matrix was scaled, then rescale eigenvalues appropriately. */
if (iscale == 1) {
if (*info == 0) {
imax = *m;
} else {
imax = *info - 1;
}
d__1 = 1. / sigma;
blasf77_sscal(&imax, &d__1, &w[1], &ione);
}
/* If eigenvalues are not in order, then sort them, along with
eigenvectors. */
if (wantz) {
for (j = 1; j <= *m-1; ++j) {
i = 0;
tmp1 = w[j];
for (jj = j + 1; jj <= *m; ++jj) {
if (w[jj] < tmp1) {
i = jj;
tmp1 = w[jj];
}
}
if (i != 0) {
itmp1 = iwork[indibl + i - 1];
w[i] = w[j];
iwork[indibl + i - 1] = iwork[indibl + j - 1];
w[j] = tmp1;
iwork[indibl + j - 1] = itmp1;
magma_cswap(n, dZ + (i-1)*lddz, ione, dZ + (j-1)*lddz, ione);
if (*info != 0) {
itmp1 = ifail[i];
ifail[i] = ifail[j];
ifail[j] = itmp1;
}
}
}
}
/* Set WORK[0] to optimal complex workspace size. */
work[1] = MAGMA_C_MAKE( lopt, 0 );
return *info;
} /* magma_cheevx_gpu */