// On input, A and ipiv is LU factorization of A. On output, A is overwritten.
// Requires m == n.
// Uses init_matrix() to re-generate original A as needed.
// Generates random RHS b and solves Ax=b.
// Returns residual, |Ax - b| / (n |A| |x|).
double get_residual(
magma_uplo_t uplo, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *A, magma_int_t lda, magma_int_t *ipiv,
magmaDoubleComplex *x, magma_int_t ldx,
magmaDoubleComplex *b, magma_int_t ldb)
{
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magma_int_t ione = 1;
// reset to original A
init_matrix( n, n, A, lda );
// compute r = Ax - b, saved in b
blasf77_zgemv( "Notrans", &n, &n, &c_one, A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
double norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_zlange( MagmaFullStr, &n, &n, A, &lda, work );
norm_r = lapackf77_zlange( MagmaFullStr, &n, &ione, b, &n, work );
norm_x = lapackf77_zlange( MagmaFullStr, &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_zprint( 1, n, b, 1 );
//printf( "r=%.2e, A=%.2e, x=%.2e, n=%d\n", norm_r, norm_A, norm_x, n );
return norm_r / (n * norm_A * norm_x);
}
// On input, A and ipiv is LU factorization of A. On output, A is overwritten.
// Requires m == n.
// Uses init_matrix() to re-generate original A as needed.
// Generates random RHS b and solves Ax=b.
// Returns residual, |Ax - b| / (n |A| |x|).
double get_residual(
magma_int_t m, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
magma_int_t *ipiv )
{
if ( m != n ) {
printf( "\nERROR: residual check defined only for square matrices\n" );
return -1;
}
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magma_int_t ione = 1;
// this seed should be DIFFERENT than used in init_matrix
// (else x is column of A, so residual can be exactly zero)
magma_int_t ISEED[4] = {0,0,0,2};
magma_int_t info = 0;
magmaDoubleComplex *x, *b;
// initialize RHS
TESTING_MALLOC_CPU( x, magmaDoubleComplex, n );
TESTING_MALLOC_CPU( b, magmaDoubleComplex, n );
lapackf77_zlarnv( &ione, ISEED, &n, b );
blasf77_zcopy( &n, b, &ione, x, &ione );
// solve Ax = b
lapackf77_zgetrs( "Notrans", &n, &ione, A, &lda, ipiv, x, &n, &info );
if (info != 0) {
printf("lapackf77_zgetrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
// reset to original A
init_matrix( m, n, A, lda );
// compute r = Ax - b, saved in b
blasf77_zgemv( "Notrans", &m, &n, &c_one, A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
double norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_zlange( "F", &m, &n, A, &lda, work );
norm_r = lapackf77_zlange( "F", &n, &ione, b, &n, work );
norm_x = lapackf77_zlange( "F", &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_zprint( 1, n, b, 1 );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
//printf( "r=%.2e, A=%.2e, x=%.2e, n=%d\n", norm_r, norm_A, norm_x, n );
return norm_r / (n * norm_A * norm_x);
}
magma_int_t magma_ztrevc3(
magma_side_t side, magma_vec_t howmany,
magma_int_t *select, // logical in Fortran
magma_int_t n,
magmaDoubleComplex *T, magma_int_t ldt,
magmaDoubleComplex *VL, magma_int_t ldvl,
magmaDoubleComplex *VR, magma_int_t ldvr,
magma_int_t mm, magma_int_t *mout,
magmaDoubleComplex *work, magma_int_t lwork,
double *rwork, magma_int_t *info )
{
#define T(i,j) ( T + (i) + (j)*ldt )
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
#define work(i,j) (work + (i) + (j)*n)
// .. Parameters ..
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magma_int_t nbmin = 16, nbmax = 128;
const magma_int_t ione = 1;
// .. Local Scalars ..
magma_int_t allv, bothv, leftv, over, rightv, somev;
magma_int_t i, ii, is, j, k, ki, iv, n2, nb, nb2, version;
double ovfl, remax, scale, smin, smlnum, ulp, unfl;
// Decode and test the input parameters
bothv = (side == MagmaBothSides);
rightv = (side == MagmaRight) || bothv;
leftv = (side == MagmaLeft ) || bothv;
allv = (howmany == MagmaAllVec);
over = (howmany == MagmaBacktransVec);
somev = (howmany == MagmaSomeVec);
// Set mout to the number of columns required to store the selected
// eigenvectors.
if ( somev ) {
*mout = 0;
for( j=0; j < n; ++j ) {
if ( select[j] ) {
*mout += 1;
}
}
}
else {
*mout = n;
}
*info = 0;
if ( ! rightv && ! leftv )
*info = -1;
else if ( ! allv && ! over && ! somev )
*info = -2;
else if ( n < 0 )
*info = -4;
else if ( ldt < max( 1, n ) )
*info = -6;
else if ( ldvl < 1 || ( leftv && ldvl < n ) )
*info = -8;
else if ( ldvr < 1 || ( rightv && ldvr < n ) )
*info = -10;
else if ( mm < *mout )
*info = -11;
else if ( lwork < max( 1, 2*n ) )
*info = -14;
if ( *info != 0 ) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// Quick return if possible.
if ( n == 0 ) {
return *info;
}
// Use blocked version (2) if sufficient workspace.
// Requires 1 vector to save diagonal elements, and 2*nb vectors for x and Q*x.
// (Compared to dtrevc3, rwork stores 1-norms.)
// Zero-out the workspace to avoid potential NaN propagation.
nb = 2;
if ( lwork >= n + 2*n*nbmin ) {
version = 2;
nb = (lwork - n) / (2*n);
nb = min( nb, nbmax );
nb2 = 1 + 2*nb;
lapackf77_zlaset( "F", &n, &nb2, &c_zero, &c_zero, work, &n );
}
else {
version = 1;
}
// Set the constants to control overflow.
unfl = lapackf77_dlamch( "Safe minimum" );
ovfl = 1. / unfl;
lapackf77_dlabad( &unfl, &ovfl );
ulp = lapackf77_dlamch( "Precision" );
smlnum = unfl*( n / ulp );
// Store the diagonal elements of T in working array work.
for( i=0; i < n; ++i ) {
*work(i,0) = *T(i,i);
}
// Compute 1-norm of each column of strictly upper triangular
// part of T to control overflow in triangular solver.
rwork[0] = 0.;
for( j=1; j < n; ++j ) {
rwork[j] = cblas_dzasum( j, T(0,j), ione );
}
magma_timer_t time_total=0, time_trsv=0, time_gemm=0, time_gemv=0, time_trsv_sum=0, time_gemm_sum=0, time_gemv_sum=0;
timer_start( time_total );
if ( rightv ) {
// ============================================================
// Compute right eigenvectors.
// iv is index of column in current block.
// Non-blocked version always uses iv=1;
// blocked version starts with iv=nb, goes down to 1.
// (Note the "0-th" column is used to store the original diagonal.)
iv = 1;
if ( version == 2 ) {
iv = nb;
}
timer_start( time_trsv );
is = *mout - 1;
for( ki=n-1; ki >= 0; --ki ) {
if ( somev ) {
if ( ! select[ki] ) {
continue;
}
}
smin = max( ulp*( MAGMA_Z_ABS1( *T(ki,ki) ) ), smlnum );
// --------------------------------------------------------
// Complex right eigenvector
*work(ki,iv) = c_one;
// Form right-hand side.
for( k=0; k < ki; ++k ) {
*work(k,iv) = -(*T(k,ki));
}
// Solve upper triangular system:
// [ T(1:ki-1,1:ki-1) - T(ki,ki) ]*X = scale*work.
for( k=0; k < ki; ++k ) {
*T(k,k) -= *T(ki,ki);
if ( MAGMA_Z_ABS1( *T(k,k) ) < smin ) {
*T(k,k) = MAGMA_Z_MAKE( smin, 0. );
}
}
if ( ki > 0 ) {
lapackf77_zlatrs( "Upper", "No transpose", "Non-unit", "Y",
&ki, T, &ldt,
work(0,iv), &scale, rwork, info );
*work(ki,iv) = MAGMA_Z_MAKE( scale, 0. );
}
// Copy the vector x or Q*x to VR and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VR and normalize
n2 = ki+1;
blasf77_zcopy( &n2, work(0,iv), &ione, VR(0,is), &ione );
ii = blasf77_izamax( &n2, VR(0,is), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *VR(ii,is) );
blasf77_zdscal( &n2, &remax, VR(0,is), &ione );
for( k=ki+1; k < n; ++k ) {
*VR(k,is) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemv );
if ( ki > 0 ) {
blasf77_zgemv( "n", &n, &ki, &c_one,
VR, &ldvr,
work(0, iv), &ione,
work(ki,iv), VR(0,ki), &ione );
}
time_gemv_sum += timer_stop( time_gemv );
ii = blasf77_izamax( &n, VR(0,ki), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *VR(ii,ki) );
blasf77_zdscal( &n, &remax, VR(0,ki), &ione );
timer_start( time_trsv );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out below vector
for( k=ki+1; k < n; ++k ) {
*work(k,iv) = c_zero;
}
// Columns iv:nb of work are valid vectors.
// When the number of vectors stored reaches nb,
// or if this was last vector, do the GEMM
if ( (iv == 1) || (ki == 0) ) {
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemm );
nb2 = nb-iv+1;
n2 = ki+nb-iv+1;
blasf77_zgemm( "n", "n", &n, &nb2, &n2, &c_one,
VR, &ldvr,
work(0,iv ), &n, &c_zero,
work(0,nb+iv), &n );
time_gemm_sum += timer_stop( time_gemm );
// normalize vectors
// TODO if somev, should copy vectors individually to correct location.
for( k = iv; k <= nb; ++k ) {
ii = blasf77_izamax( &n, work(0,nb+k), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *work(ii,nb+k) );
blasf77_zdscal( &n, &remax, work(0,nb+k), &ione );
}
lapackf77_zlacpy( "F", &n, &nb2, work(0,nb+iv), &n, VR(0,ki), &ldvr );
iv = nb;
timer_start( time_trsv );
}
else {
iv -= 1;
}
} // blocked back-transform
// Restore the original diagonal elements of T.
for( k=0; k <= ki - 1; ++k ) {
*T(k,k) = *work(k,0);
}
is -= 1;
}
}
timer_stop( time_trsv );
timer_stop( time_total );
timer_printf( "trevc trsv %.4f, gemm %.4f, gemv %.4f, total %.4f\n",
time_trsv_sum, time_gemm_sum, time_gemv_sum, time_total );
if ( leftv ) {
// ============================================================
// Compute left eigenvectors.
// iv is index of column in current block.
// Non-blocked version always uses iv=1;
// blocked version starts with iv=1, goes up to nb.
// (Note the "0-th" column is used to store the original diagonal.)
iv = 1;
is = 0;
for( ki=0; ki < n; ++ki ) {
if ( somev ) {
if ( ! select[ki] ) {
continue;
}
}
smin = max( ulp*MAGMA_Z_ABS1( *T(ki,ki) ), smlnum );
// --------------------------------------------------------
// Complex left eigenvector
*work(ki,iv) = c_one;
// Form right-hand side.
for( k = ki + 1; k < n; ++k ) {
*work(k,iv) = -MAGMA_Z_CNJG( *T(ki,k) );
}
// Solve conjugate-transposed triangular system:
// [ T(ki+1:n,ki+1:n) - T(ki,ki) ]**H * X = scale*work.
for( k = ki + 1; k < n; ++k ) {
*T(k,k) -= *T(ki,ki);
if ( MAGMA_Z_ABS1( *T(k,k) ) < smin ) {
*T(k,k) = MAGMA_Z_MAKE( smin, 0. );
}
}
if ( ki < n-1 ) {
n2 = n-ki-1;
lapackf77_zlatrs( "Upper", "Conjugate transpose", "Non-unit", "Y",
&n2, T(ki+1,ki+1), &ldt,
work(ki+1,iv), &scale, rwork, info );
*work(ki,iv) = MAGMA_Z_MAKE( scale, 0. );
}
// Copy the vector x or Q*x to VL and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VL and normalize
n2 = n-ki;
blasf77_zcopy( &n2, work(ki,iv), &ione, VL(ki,is), &ione );
ii = blasf77_izamax( &n2, VL(ki,is), &ione ) + ki - 1;
remax = 1. / MAGMA_Z_ABS1( *VL(ii,is) );
blasf77_zdscal( &n2, &remax, VL(ki,is), &ione );
for( k=0; k < ki; ++k ) {
*VL(k,is) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
if ( ki < n-1 ) {
n2 = n-ki-1;
blasf77_zgemv( "n", &n, &n2, &c_one,
VL(0,ki+1), &ldvl,
work(ki+1,iv), &ione,
work(ki, iv), VL(0,ki), &ione );
}
ii = blasf77_izamax( &n, VL(0,ki), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *VL(ii,ki) );
blasf77_zdscal( &n, &remax, VL(0,ki), &ione );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out above vector
// could go from (ki+1)-NV+1 to ki
for( k=0; k < ki; ++k ) {
*work(k,iv) = c_zero;
}
// Columns 1:iv of work are valid vectors.
// When the number of vectors stored reaches nb,
// or if this was last vector, do the GEMM
if ( (iv == nb) || (ki == n-1) ) {
n2 = n-(ki+1)+iv;
blasf77_zgemm( "n", "n", &n, &iv, &n2, &c_one,
VL(0,ki-iv+1), &ldvl,
work(ki-iv+1,1 ), &n, &c_zero,
work(0, nb+1), &n );
// normalize vectors
for( k=1; k <= iv; ++k ) {
ii = blasf77_izamax( &n, work(0,nb+k), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *work(ii,nb+k) );
blasf77_zdscal( &n, &remax, work(0,nb+k), &ione );
}
lapackf77_zlacpy( "F", &n, &iv, work(0,nb+1), &n, VL(0,ki-iv+1), &ldvl );
iv = 1;
}
else {
iv += 1;
}
} // blocked back-transform
// Restore the original diagonal elements of T.
for( k = ki + 1; k < n; ++k ) {
*T(k,k) = *work(k,0);
}
is += 1;
}
}
return *info;
} // End of ZTREVC
/**
Purpose
-------
ZLAQPS computes a step of QR factorization with column pivoting
of a complex M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
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]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau COMPLEX_16 array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 DOUBLE PRECISION array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 DOUBLE PRECISION array, dimension (N)
The vector with the exact column norms.
@param[in,out]
auxv COMPLEX_16 array, dimension (NB)
Auxiliar vector.
@param[in,out]
F COMPLEX_16 array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
@param[in]
ldf INTEGER
The leading dimension of the array F. LDF >= max(1,N).
@ingroup magma_zgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_zlaqps(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaDoubleComplex *tau, double *vn1, double *vn2,
magmaDoubleComplex *auxv,
magmaDoubleComplex *F, magma_int_t ldf,
magmaDoubleComplex_ptr dF, magma_int_t lddf)
{
#define A(i, j) (A + (i) + (j)*(lda ))
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define F(i, j) (F + (i) + (j)*(ldf ))
#define dF(i, j) (dF + (i) + (j)*(lddf))
magmaDoubleComplex c_zero = MAGMA_Z_MAKE( 0.,0.);
magmaDoubleComplex c_one = MAGMA_Z_MAKE( 1.,0.);
magmaDoubleComplex c_neg_one = MAGMA_Z_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
double d__1;
magmaDoubleComplex z__1;
magma_int_t j, k, rk;
magmaDoubleComplex Akk;
magma_int_t pvt;
double temp, temp2, tol3z;
magma_int_t itemp;
magma_int_t lsticc;
magma_int_t lastrk;
lastrk = min( m, n + offset );
tol3z = magma_dsqrt( lapackf77_dlamch("Epsilon"));
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
lsticc = 0;
k = 0;
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran idamax; pvt, k are 0-based.
i__1 = n-k;
pvt = k + blasf77_idamax( &i__1, &vn1[k], &ione ) - 1;
if (pvt != k) {
if (pvt >= nb) {
/* 1. Start copy from GPU */
magma_zgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, queue );
}
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
blasf77_zswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
vn1[pvt] = vn1[k];
vn2[pvt] = vn2[k];
if (pvt < nb) {
/* no need of transfer if pivot is within the panel */
blasf77_zswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
/* 1. Finish copy from GPU */
magma_queue_sync( queue );
/* 2. Swap as usual on CPU */
blasf77_zswap(&m, A(0, pvt), &ione, A(0, k), &ione);
/* 3. Restore the GPU */
magma_zsetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, queue );
}
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
#ifdef COMPLEX
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_Z_CONJ( *F(k,j) );
}
#endif
i__1 = m - rk;
i__2 = k;
blasf77_zgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );
#ifdef COMPLEX
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_Z_CONJ( *F(k,j) );
}
#endif
}
/* Generate elementary reflector H(k). */
if (rk < m-1) {
i__1 = m - rk;
lapackf77_zlarfg( &i__1, A(rk, k), A(rk + 1, k), &ione, &tau[k] );
} else {
lapackf77_zlarfg( &ione, A(rk, k), A(rk, k), &ione, &tau[k] );
}
Akk = *A(rk, k);
*A(rk, k) = c_one;
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
magma_zsetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda, queue );
/* Multiply on GPU */
// was CALL ZGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
magma_int_t i__3 = nb-k-1;
magma_int_t i__4 = i__2 - i__3;
magma_int_t i__5 = nb-k;
magma_zgemv( MagmaConjTrans, i__1 - i__5, i__2 - i__3,
tau[k], dA(rk +i__5, k+1+i__3), ldda,
dA(rk +i__5, k ), ione,
c_zero, dF(k+1+i__3, k ), ione, queue );
magma_zgetmatrix_async( i__2-i__3, 1,
dF(k + 1 +i__3, k), i__2,
F (k + 1 +i__3, k), i__2, queue );
blasf77_zgemv( MagmaConjTransStr, &i__1, &i__3,
&tau[k], A(rk, k+1), &lda,
A(rk, k ), &ione,
&c_zero, F(k+1, k ), &ione );
magma_queue_sync( queue );
blasf77_zgemv( MagmaConjTransStr, &i__5, &i__4,
&tau[k], A(rk, k+1+i__3), &lda,
A(rk, k ), &ione,
&c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros. */
for (j = 0; j < k; ++j) {
*F(j, k) = c_zero;
}
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K). */
if (k > 0) {
i__1 = m - rk;
i__2 = k;
z__1 = MAGMA_Z_NEGATE( tau[k] );
blasf77_zgemv( MagmaConjTransStr, &i__1, &i__2,
&z__1, A(rk, 0), &lda,
A(rk, k), &ione,
&c_zero, auxv, &ione );
i__1 = k;
blasf77_zgemv( MagmaNoTransStr, &n, &i__1,
&c_one, F(0,0), &ldf,
auxv, &ione,
&c_one, F(0,k), &ione );
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
blasf77_zgemm( MagmaNoTransStr, MagmaConjTransStr, &ione, &i__1, &i__2,
&c_neg_one, A(rk, 0 ), &lda,
F(k+1,0 ), &ldf,
&c_one, A(rk, k+1), &lda );
}
/* Update partial column norms. */
if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
/* NOTE: The following 4 lines follow from the analysis in
Lapack Working Note 176. */
temp = MAGMA_Z_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (double) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_dsqrt(temp);
}
}
}
}
*A(rk, k) = Akk;
++k;
}
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU */
magma_zsetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2, queue );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), i__2,
c_one, dA(rk+1, *kb), ldda, queue );
}
/* Recomputation of difficult columns. */
while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb) {
vn1[lsticc] = magma_cblas_dznrm2( i__1, A(rk+1,lsticc), ione );
}
else {
/* Where is the data, CPU or GPU ? */
double r1, r2;
r1 = magma_cblas_dznrm2( nb-k, A(rk+1,lsticc), ione );
r2 = magma_dznrm2( m-offset-nb, dA(offset + nb + 1, lsticc), ione, queue );
//vn1[lsticc] = magma_dznrm2( i__1, dA(rk + 1, lsticc), ione, queue );
vn1[lsticc] = magma_dsqrt(r1*r1 + r2*r2);
}
/* NOTE: The computation of VN1( LSTICC ) relies on the fact that
SNRM2 does not fail on vectors with norm below the value of SQRT(DLAMCH('S')) */
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;
}
magma_queue_destroy( queue );
return MAGMA_SUCCESS;
} /* magma_zlaqps */
extern "C" magma_int_t
magma_zlabrd_gpu(
magma_int_t m, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex_ptr da, size_t da_offset, magma_int_t ldda,
double *d, double *e, magmaDoubleComplex *tauq, magmaDoubleComplex *taup,
magmaDoubleComplex *x, magma_int_t ldx,
magmaDoubleComplex_ptr dx, size_t dx_offset, magma_int_t lddx,
magmaDoubleComplex *y, magma_int_t ldy,
magmaDoubleComplex_ptr dy, size_t dy_offset, magma_int_t lddy,
magma_queue_t queue )
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
ZLABRD reduces the first NB rows and columns of a complex general
m by n matrix A to upper or lower bidiagonal form by an orthogonal
transformation Q' * A * P, and returns the matrices X and Y which
are needed to apply the transformation to the unreduced part of A.
If m >= n, A is reduced to upper bidiagonal form; if m < n, to lower
bidiagonal form.
This is an auxiliary routine called by SGEBRD
Arguments
=========
M (input) INTEGER
The number of rows in the matrix A.
N (input) INTEGER
The number of columns in the matrix A.
NB (input) INTEGER
The number of leading rows and columns of A to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the m by n general matrix to be reduced.
On exit, the first NB rows and columns of the matrix are
overwritten; the rest of the array is unchanged.
If m >= n, elements on and below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors; and
elements above the diagonal in the first NB rows, with the
array TAUP, represent the orthogonal matrix P as a product
of elementary reflectors.
If m < n, elements below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors, and
elements on and above the diagonal in the first NB rows,
with the array TAUP, represent the orthogonal matrix P as
a product of elementary reflectors.
See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
D (output) COMPLEX_16 array, dimension (NB)
The diagonal elements of the first NB rows and columns of
the reduced matrix. D(i) = A(i,i).
E (output) COMPLEX_16 array, dimension (NB)
The off-diagonal elements of the first NB rows and columns of
the reduced matrix.
TAUQ (output) COMPLEX_16 array dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
TAUP (output) COMPLEX_16 array, dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
X (output) COMPLEX_16 array, dimension (LDX,NB)
The m-by-nb matrix X required to update the unreduced part
of A.
LDX (input) INTEGER
The leading dimension of the array X. LDX >= M.
Y (output) COMPLEX_16 array, dimension (LDY,NB)
The n-by-nb matrix Y required to update the unreduced part
of A.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
Further Details
===============
The matrices Q and P are represented as products of elementary
reflectors:
Q = H(1) H(2) . . . H(nb) and P = G(1) G(2) . . . G(nb)
Each H(i) and G(i) has the form:
H(i) = I - tauq * v * v' and G(i) = I - taup * u * u'
where tauq and taup are complex scalars, and v and u are complex vectors.
If m >= n, v(1:i-1) = 0, v(i) = 1, and v(i:m) is stored on exit in
A(i:m,i); u(1:i) = 0, u(i+1) = 1, and u(i+1:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
If m < n, v(1:i) = 0, v(i+1) = 1, and v(i+1:m) is stored on exit in
A(i+2:m,i); u(1:i-1) = 0, u(i) = 1, and u(i:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
The elements of the vectors v and u together form the m-by-nb matrix
V and the nb-by-n matrix U' which are needed, with X and Y, to apply
the transformation to the unreduced part of the matrix, using a block
update of the form: A := A - V*Y' - X*U'.
The contents of A on exit are illustrated by the following examples
with nb = 2:
m = 6 and n = 5 (m > n): m = 5 and n = 6 (m < n):
( 1 1 u1 u1 u1 ) ( 1 u1 u1 u1 u1 u1 )
( v1 1 1 u2 u2 ) ( 1 1 u2 u2 u2 u2 )
( v1 v2 a a a ) ( v1 1 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a )
where a denotes an element of the original matrix which is unchanged,
vi denotes an element of the vector defining H(i), and ui an element
of the vector defining G(i).
===================================================================== */
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magma_int_t c__1 = 1;
magma_int_t a_dim1, a_offset, x_dim1, x_offset, y_dim1, y_offset, i__2, i__3;
magma_int_t i__;
magmaDoubleComplex alpha;
a_dim1 = lda;
a_offset = 1 + a_dim1;
a -= a_offset;
--d;
--e;
--tauq;
--taup;
x_dim1 = ldx;
x_offset = 1 + x_dim1;
x -= x_offset;
dx_offset -= 1 + lddx;
y_dim1 = ldy;
y_offset = 1 + y_dim1;
y -= y_offset;
dy_offset -= 1 + lddy;
/* Quick return if possible */
if (m <= 0 || n <= 0) {
return 0;
}
magmaDoubleComplex *f;
magma_zmalloc_cpu( &f, max(n,m) );
assert( f != NULL ); // TODO return error, or allocate outside zlatrd
magma_event_t event = NULL;
if (m >= n) {
/* Reduce to upper bidiagonal form */
for (i__ = 1; i__ <= nb; ++i__) {
/* Update A(i:m,i) */
i__2 = m - i__ + 1;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i__3, &y[i__+y_dim1], &ldy );
#endif
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one, &a[i__ + a_dim1], &lda,
&y[i__+y_dim1], &ldy, &c_one, &a[i__ + i__ * a_dim1], &c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i__3, &y[i__+y_dim1], &ldy );
#endif
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one, &x[i__ + x_dim1], &ldx,
&a[i__*a_dim1+1], &c__1, &c_one, &a[i__+i__*a_dim1], &c__1);
/* Generate reflection Q(i) to annihilate A(i+1:m,i) */
alpha = a[i__ + i__ * a_dim1];
i__2 = m - i__ + 1;
i__3 = i__ + 1;
lapackf77_zlarfg(&i__2, &alpha,
&a[min(i__3,m) + i__ * a_dim1], &c__1, &tauq[i__]);
d[i__] = MAGMA_Z_REAL( alpha );
if (i__ < n) {
a[i__ + i__ * a_dim1] = c_one;
/* Compute Y(i+1:n,i) */
i__2 = m - i__ + 1;
i__3 = n - i__;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_zsetvector( i__2,
a + i__ + i__ * a_dim1, 1,
da, da_offset + (i__-1)+(i__-1)* (ldda), 1,
queue );
// 2. Multiply ---------------------------------------------
magma_zgemv(MagmaConjTrans, i__2, i__3, c_one,
da, da_offset + (i__-1) + ((i__-1) + 1) * (ldda), ldda,
da, da_offset + (i__-1) + (i__-1) * (ldda), c__1, c_zero,
dy, dy_offset + i__ + 1 + i__ * y_dim1, c__1,
queue );
// 3. Put the result back ----------------------------------
magma_zgetmatrix_async( i__3, 1,
dy, dy_offset + i__+1+i__*y_dim1, y_dim1,
y+i__+1+i__*y_dim1, y_dim1,
queue, &event );
i__2 = m - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one, &a[i__ + a_dim1],
&lda, &a[i__ + i__ * a_dim1], &c__1, &c_zero,
&y[i__ * y_dim1 + 1], &c__1);
i__2 = n - i__;
i__3 = i__ - 1;
blasf77_zgemv("N", &i__2, &i__3, &c_neg_one, &y[i__ + 1 +y_dim1], &ldy,
&y[i__ * y_dim1 + 1], &c__1,
&c_zero, f, &c__1);
i__2 = m - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one, &x[i__ + x_dim1],
&ldx, &a[i__ + i__ * a_dim1], &c__1, &c_zero,
&y[i__ * y_dim1 + 1], &c__1);
// 4. Synch to make sure the result is back ----------------
magma_event_sync( event );
if (i__3 != 0){
i__2 = n - i__;
blasf77_zaxpy(&i__2, &c_one, f,&c__1, &y[i__+1+i__*y_dim1],&c__1);
}
i__2 = i__ - 1;
i__3 = n - i__;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
&a[(i__ + 1) * a_dim1 + 1], &lda, &y[i__ * y_dim1 + 1], &c__1, &c_one,
&y[i__ + 1 + i__ * y_dim1], &c__1);
i__2 = n - i__;
blasf77_zscal(&i__2, &tauq[i__], &y[i__ + 1 + i__ * y_dim1], &c__1);
/* Update A(i,i+1:n) */
i__2 = n - i__;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i__2, &a[i__+(i__+1)*a_dim1], &lda );
lapackf77_zlacgv( &i__, &a[i__+a_dim1], &lda );
#endif
blasf77_zgemv("No transpose", &i__2, &i__, &c_neg_one,
&y[i__ + 1 + y_dim1], &ldy, &a[i__ + a_dim1], &lda,
&c_one, &a[i__ + (i__ + 1) * a_dim1], &lda);
i__2 = i__ - 1;
i__3 = n - i__;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i__, &a[i__+a_dim1], &lda );
lapackf77_zlacgv( &i__2, &x[i__+x_dim1], &ldx );
#endif
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_neg_one, &a[(i__ + 1) *
a_dim1 + 1], &lda, &x[i__ + x_dim1], &ldx, &c_one, &a[
i__ + (i__ + 1) * a_dim1], &lda);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i__2, &x[i__+x_dim1], &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+2:n) */
i__2 = n - i__;
/* Computing MIN */
i__3 = i__ + 2;
alpha = a[i__ + (i__ + 1) * a_dim1];
lapackf77_zlarfg(&i__2, &alpha, &a[i__ + min(
i__3,n) * a_dim1], &lda, &taup[i__]);
e[i__] = MAGMA_Z_REAL( alpha );
a[i__ + (i__ + 1) * a_dim1] = c_one;
/* Compute X(i+1:m,i) */
i__2 = m - i__;
i__3 = n - i__;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_zsetvector( i__3,
a + i__ + (i__ +1)* a_dim1, lda,
da, da_offset + (i__-1)+((i__-1)+1)*(ldda), ldda,
queue );
// 2. Multiply ---------------------------------------------
//magma_zcopy(i__3, da+(i__-1)+((i__-1)+1)*(ldda), ldda,
// dy + 1 + lddy, 1);
magma_zgemv(MagmaNoTrans, i__2, i__3, c_one,
da, da_offset + (i__-1)+1+ ((i__-1)+1) * (ldda), ldda,
da, da_offset + (i__-1) + ((i__-1)+1) * (ldda), ldda,
//dy + 1 + lddy, 1,
c_zero, dx, dx_offset + i__ + 1 + i__ * x_dim1, c__1,
queue );
// 3. Put the result back ----------------------------------
magma_zgetmatrix_async( i__2, 1,
dx, dx_offset + i__+1+i__*x_dim1, x_dim1,
x+i__+1+i__*x_dim1, x_dim1,
queue, &event );
i__2 = n - i__;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__, &c_one, &y[i__ + 1 + y_dim1],
&ldy, &a[i__ + (i__ + 1) * a_dim1], &lda, &c_zero, &x[
i__ * x_dim1 + 1], &c__1);
i__2 = m - i__;
blasf77_zgemv("N", &i__2, &i__, &c_neg_one, &a[i__ + 1 + a_dim1], &lda,
&x[i__ * x_dim1 + 1], &c__1, &c_zero, f, &c__1);
i__2 = i__ - 1;
i__3 = n - i__;
blasf77_zgemv("N", &i__2, &i__3, &c_one, &a[(i__ + 1) * a_dim1 + 1],
&lda, &a[i__ + (i__ + 1) * a_dim1], &lda,
&c_zero, &x[i__ * x_dim1 + 1], &c__1);
// 4. Synch to make sure the result is back ----------------
magma_event_sync( event );
if (i__!=0){
i__2 = m - i__;
blasf77_zaxpy(&i__2, &c_one, f,&c__1, &x[i__+1+i__*x_dim1],&c__1);
}
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one, &x[i__ + 1 +
x_dim1], &ldx, &x[i__ * x_dim1 + 1], &c__1, &c_one, &x[
i__ + 1 + i__ * x_dim1], &c__1);
i__2 = m - i__;
blasf77_zscal(&i__2, &taup[i__], &x[i__ + 1 + i__ * x_dim1], &c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
i__2 = n - i__;
lapackf77_zlacgv( &i__2, &a[i__+(i__+1)*a_dim1], &lda );
// 4. Send the block reflector A(i+1:m,i) to the GPU after ZLACGV()
magma_zsetvector( i__2,
a + i__ + (i__ +1)* a_dim1, lda,
da, da_offset + (i__-1)+((i__-1)+1)*(ldda), ldda,
queue );
#endif
}
}
}
else {
/* Reduce to lower bidiagonal form */
for (i__ = 1; i__ <= nb; ++i__) {
/* Update A(i,i:n) */
i__2 = n - i__ + 1;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__2, &a[i__ + i__ * a_dim1], &lda);
lapackf77_zlacgv(&i__3, &a[i__ + a_dim1], &lda);
#endif
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one, &y[i__ + y_dim1], &ldy,
&a[i__ + a_dim1], &lda, &c_one, &a[i__ + i__ * a_dim1], &lda);
i__2 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &a[i__ + a_dim1], &lda);
lapackf77_zlacgv(&i__3, &x[i__ + x_dim1], &ldx);
#endif
i__3 = n - i__ + 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_neg_one, &a[i__ * a_dim1 + 1],
&lda, &x[i__ + x_dim1], &ldx, &c_one, &a[i__ + i__ * a_dim1], &lda);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__2, &x[i__ + x_dim1], &ldx);
#endif
/* Generate reflection P(i) to annihilate A(i,i+1:n) */
i__2 = n - i__ + 1;
/* Computing MIN */
i__3 = i__ + 1;
alpha = a[i__ + i__ * a_dim1];
lapackf77_zlarfg(&i__2, &alpha,
&a[i__ + min(i__3,n) * a_dim1], &lda, &taup[i__]);
d[i__] = MAGMA_Z_REAL( alpha );
if (i__ < m) {
a[i__ + i__ * a_dim1] = c_one;
/* Compute X(i+1:m,i) */
i__2 = m - i__;
i__3 = n - i__ + 1;
// 1. Send the block reflector A(i,i+1:n) to the GPU ------
magma_zsetvector( i__3,
a + i__ + i__ * a_dim1, lda,
da, da_offset + (i__-1)+(i__-1)* (ldda), ldda,
queue );
// 2. Multiply ---------------------------------------------
//magma_zcopy(i__3, da+(i__-1)+(i__-1)*(ldda), ldda,
// dy + 1 + lddy, 1);
magma_zgemv(MagmaNoTrans, i__2, i__3, c_one,
da, da_offset + (i__-1)+1 + (i__-1) * ldda, ldda,
da, da_offset + (i__-1) + (i__-1) * ldda, ldda,
// dy + 1 + lddy, 1,
c_zero,
dx, dx_offset + i__ + 1 + i__ * x_dim1, c__1,
queue );
// 3. Put the result back ----------------------------------
magma_zgetmatrix_async( i__2, 1,
dx, dx_offset + i__+1+i__*x_dim1, x_dim1,
x+i__+1+i__*x_dim1, x_dim1,
queue, &event );
i__2 = n - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one, &y[i__ + y_dim1],
&ldy, &a[i__ + i__ * a_dim1], &lda, &c_zero,
&x[i__ * x_dim1 + 1], &c__1);
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one,
&a[i__ + 1 + a_dim1], &lda, &x[i__ * x_dim1 + 1], &c__1, &c_zero,
f, &c__1);
i__2 = i__ - 1;
i__3 = n - i__ + 1;
blasf77_zgemv("No transpose", &i__2, &i__3, &c_one,
&a[i__ * a_dim1 + 1], &lda, &a[i__ + i__ * a_dim1], &lda, &c_zero,
&x[i__ * x_dim1 + 1], &c__1);
// 4. Synch to make sure the result is back ----------------
magma_event_sync( event );
if (i__2 != 0){
i__3 = m - i__;
blasf77_zaxpy(&i__3, &c_one, f,&c__1, &x[i__+1+i__*x_dim1],&c__1);
}
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one,
&x[i__ + 1 + x_dim1], &ldx, &x[i__ * x_dim1 + 1], &c__1, &c_one,
&x[i__ + 1 + i__ * x_dim1], &c__1);
i__2 = m - i__;
blasf77_zscal(&i__2, &taup[i__], &x[i__ + 1 + i__ * x_dim1], &c__1);
i__2 = n - i__ + 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__2, &a[i__ + i__ * a_dim1], &lda);
magma_zsetvector( i__2,
a + i__ + (i__ )* a_dim1, lda,
da, da_offset + (i__-1)+ (i__-1)*(ldda), ldda,
queue );
#endif
/* Update A(i+1:m,i) */
i__2 = m - i__;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &y[i__ + y_dim1], &ldy);
#endif
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one,
&a[i__ + 1 + a_dim1], &lda, &y[i__ + y_dim1], &ldy, &c_one,
&a[i__ + 1 + i__ * a_dim1], &c__1);
i__2 = m - i__;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &y[i__ + y_dim1], &ldy);
#endif
blasf77_zgemv("No transpose", &i__2, &i__, &c_neg_one,
&x[i__ + 1 + x_dim1], &ldx, &a[i__ * a_dim1 + 1], &c__1, &c_one,
&a[i__ + 1 + i__ * a_dim1], &c__1);
/* Generate reflection Q(i) to annihilate A(i+2:m,i) */
i__2 = m - i__;
i__3 = i__ + 2;
alpha = a[i__ + 1 + i__ * a_dim1];
lapackf77_zlarfg(&i__2, &alpha,
&a[min(i__3,m) + i__ * a_dim1], &c__1, &tauq[i__]);
e[i__] = MAGMA_Z_REAL( alpha );
a[i__ + 1 + i__ * a_dim1] = c_one;
/* Compute Y(i+1:n,i) */
i__2 = m - i__;
i__3 = n - i__;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_zsetvector( i__2,
a + i__ +1+ i__ * a_dim1, 1,
da, da_offset + (i__-1)+1+ (i__-1)*(ldda), 1,
queue );
// 2. Multiply ---------------------------------------------
magma_zgemv(MagmaConjTrans, i__2, i__3, c_one,
da, da_offset + (i__-1)+1+ ((i__-1)+1) * ldda, ldda,
da, da_offset + (i__-1)+1+ (i__-1) * ldda, c__1,
c_zero, dy, dy_offset + i__ + 1 + i__ * y_dim1, c__1,
queue );
// 3. Put the result back ----------------------------------
magma_zgetmatrix_async( i__3, 1,
dy, dy_offset + i__+1+i__*y_dim1, y_dim1,
y+i__+1+i__*y_dim1, y_dim1,
queue, &event );
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one, &a[i__ + 1 + a_dim1],
&lda, &a[i__ + 1 + i__ * a_dim1], &c__1, &c_zero,
&y[ i__ * y_dim1 + 1], &c__1);
i__2 = n - i__;
i__3 = i__ - 1;
blasf77_zgemv("No transpose", &i__2, &i__3, &c_neg_one,
&y[i__ + 1 + y_dim1], &ldy, &y[i__ * y_dim1 + 1], &c__1,
&c_zero, f, &c__1);
i__2 = m - i__;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__, &c_one, &x[i__ + 1 + x_dim1],
&ldx, &a[i__ + 1 + i__ * a_dim1], &c__1, &c_zero,
&y[i__ * y_dim1 + 1], &c__1);
// 4. Synch to make sure the result is back ----------------
magma_event_sync( event );
if (i__3 != 0){
i__2 = n - i__;
blasf77_zaxpy(&i__2, &c_one, f,&c__1, &y[i__+1+i__*y_dim1],&c__1);
}
i__2 = n - i__;
blasf77_zgemv(MagmaConjTransStr, &i__, &i__2, &c_neg_one,
&a[(i__ + 1) * a_dim1 + 1], &lda, &y[i__ * y_dim1 + 1],
&c__1, &c_one, &y[i__ + 1 + i__ * y_dim1], &c__1);
i__2 = n - i__;
blasf77_zscal(&i__2, &tauq[i__], &y[i__ + 1 + i__ * y_dim1], &c__1);
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else {
i__2 = n - i__ + 1;
lapackf77_zlacgv(&i__2, &a[i__ + i__ * a_dim1], &lda);
magma_zsetvector( i__2,
a + i__ + (i__ )* a_dim1, lda,
da, da_offset + (i__-1)+ (i__-1)*(ldda), ldda,
queue );
}
#endif
}
}
magma_queue_sync( queue );
magma_free_cpu(f);
return MAGMA_SUCCESS;
} /* magma_zlabrd */
int main(int argc, char **argv)
{
TESTING_CUDA_INIT();
magma_timestr_t start, end;
double flops, magma_perf, cuda_perf, error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
FILE *fp ;
magma_int_t i, lda, Xm, Ym;
magma_int_t M, M0 = 0;
magma_int_t N, N0 = 0;
magma_int_t szeA, szeX, szeY;
magma_int_t istart = 64;
magma_int_t iend = 10240;
magma_int_t incx = 1;
magma_int_t incy = 1;
char trans = MagmaNoTrans;
cuDoubleComplex alpha = MAGMA_Z_MAKE(1., 0.); // MAGMA_Z_MAKE( 1.5, -2.3 );
cuDoubleComplex beta = MAGMA_Z_MAKE(0., 0.); // MAGMA_Z_MAKE( -0.6, 0.8 );
cuDoubleComplex *A, *X, *Y, *Ycublas, *Ymagma;
cuDoubleComplex *dA, *dX, *dY;
if (argc != 1){
for(i=1; i<argc; i++){
if ( strcmp("-n", argv[i]) == 0 ){
N0 = atoi(argv[++i]);
}
else if ( strcmp("-m", argv[i]) == 0 ){
M0 = atoi(argv[++i]);
}
else if (strcmp("-N", argv[i])==0){
trans = MagmaNoTrans;
}
else if (strcmp("-T", argv[i])==0){
trans = MagmaTrans;
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else if (strcmp("-C", argv[i])==0){
trans = MagmaConjTrans;
}
#endif
}
}
if ( (M0 != 0) && (N0 != 0) )
iend = istart + 1;
M = N = iend;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if( trans == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
lda = ((M+31)/32)*32;
szeA = lda*N;
szeX = incx*Xm;
szeY = incy*Ym;
TESTING_MALLOC( A, cuDoubleComplex, szeA );
TESTING_MALLOC( X, cuDoubleComplex, szeX );
TESTING_MALLOC( Y, cuDoubleComplex, szeY );
TESTING_MALLOC( Ycublas, cuDoubleComplex, szeY );
TESTING_MALLOC( Ymagma, cuDoubleComplex, szeY );
TESTING_DEVALLOC( dA, cuDoubleComplex, szeA );
TESTING_DEVALLOC( dX, cuDoubleComplex, szeX );
TESTING_DEVALLOC( dY, cuDoubleComplex, szeY );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &szeA, A );
lapackf77_zlarnv( &ione, ISEED, &szeX, X );
lapackf77_zlarnv( &ione, ISEED, &szeY, Y );
fp = fopen ("results_zgemv.txt", "w") ;
if( fp == NULL ){ printf("Couldn't open output file\n"); exit(1);}
printf("\nUsage: \n");
printf(" testing_zgemv [-N|T|C] [-m %d] [-n %d]\n\n", 1024, 1024);
printf( " m n CUBLAS,Gflop/s MAGMABLAS Gflop/s \"error\"\n"
"==============================================================\n");
fprintf(fp, " m n CUBLAS,Gflop/s MAGMABLAS Gflop/s \"error\"\n"
"==============================================================\n");
for( i=istart; i < iend; i = (int)((i+1)*1.1) )
{
M = N = i;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if( trans == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
lda = ((M+31)/32)*32;
flops = FLOPS( (double)M, (double)N ) / 1000000;
printf( "%5d %5d ", (int) M, (int) N );
fprintf( fp, "%5d %5d ", (int) M, (int) N );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( M, N, A, lda, dA, lda );
magma_zsetvector( Xm, X, incx, dX, incx );
magma_zsetvector( Ym, Y, incy, dY, incy );
/*
* Cublas Version
*/
start = get_current_time();
cublasZgemv( trans, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
end = get_current_time();
magma_zgetvector( Ym, dY, incy, Ycublas, incy );
cuda_perf = flops / GetTimerValue(start, end);
printf( "%11.2f", cuda_perf );
fprintf(fp, "%11.2f", cuda_perf );
/*
* Magma Version
*/
magma_zsetvector( Ym, Y, incy, dY, incy );
start = get_current_time();
magmablas_zgemv( trans, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
end = get_current_time();
magma_zgetvector( Ym, dY, incx, Ymagma, incx );
magma_perf = flops / GetTimerValue(start, end);
printf( "%11.2f", magma_perf );
fprintf(fp, "%11.2f", magma_perf );
/* =====================================================================
Computing the Difference Cublas VS Magma
=================================================================== */
blasf77_zaxpy( &Ym, &c_neg_one, Ymagma, &incy, Ycublas, &incy);
error = lapackf77_zlange( "M", &Ym, &ione, Ycublas, &Ym, work );
#if 0
printf( "\t\t %8.6e", error / (double)Ym );
fprintf( fp, "\t\t %8.6e", error / (double)Ym );
/*
* Blas comparaison
*/
{
char *blastrans = MagmaNoTransStr;
if ( trans == MagmaConjTrans )
blastrans = MagmaConjTransStr;
else if ( trans == MagmaTrans )
blastrans = MagmaTransStr;
blasf77_zcopy( &Ym, Y, &incy, Ycublas, &incy );
blasf77_zgemv( blastrans, &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Ycublas, &incy );
blasf77_zaxpy( &Ym, &c_neg_one, Ymagma, &incy, Ycublas, &incy);
error = lapackf77_zlange( "M", &Ym, &ione, Ycublas, &Ym, work );
}
#endif
printf( "\t\t %8.6e\n", error / (double)Ym );
fprintf( fp, "\t\t %8.6e\n", error / (double)Ym );
}
/* Free Memory */
TESTING_FREE( A );
TESTING_FREE( X );
TESTING_FREE( Y );
TESTING_FREE( Ycublas );
TESTING_FREE( Ymagma );
TESTING_DEVFREE( dA );
TESTING_DEVFREE( dX );
TESTING_DEVFREE( dY );
/* Free device */
TESTING_CUDA_FINALIZE();
return EXIT_SUCCESS;
}
extern "C" magma_int_t
magma_zlahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *tau,
magmaDoubleComplex *T, magma_int_t ldt,
magmaDoubleComplex *Y, magma_int_t ldy,
struct zgehrd_data* data )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZLAHR2 reduces the first NB columns of a complex general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by ZGEHRD.
Arguments
=========
N (input) INTEGER
The order of the matrix A.
K (input) INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
NB (input) INTEGER
The number of columns to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX_16 array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) COMPLEX_16 array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) COMPLEX_16 array, dimension (LDY,NB)
The n-by-nb matrix Y.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
dA (input/output) COMPLEX_16 array on the GPU, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements in rows K:N of the first NB columns are
overwritten with the matrix Y.
DV (output) COMPLEX_16 array on the GPU, dimension (N, NB)
On exit this contains the Householder vectors of the transformation.
Further Details
===============
The matrix Q is represented as a product of nb 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+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
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.
===================================================================== */
#define A( i, j ) ( A + (i) + (j)*lda)
#define Y( i, j ) ( Y + (i) + (j)*ldy)
#define T( i, j ) ( T + (i) + (j)*ldt)
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex tmp;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t ione = 1;
magma_int_t d, dki1, dn, nblocks, gblock, lblock, lgid;
magma_int_t n_k_i_1, n_k;
magmaDoubleComplex scale;
magma_int_t i;
magmaDoubleComplex ei = MAGMA_Z_ZERO;
magma_int_t info_data = 0;
magma_int_t *info = &info_data;
if (n < 0) {
*info = -1;
} else if (k < 0 || k >= n) {
*info = -2;
} else if (nb < 1 || nb > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldt < nb) {
*info = -8;
} else if (ldy < max(1,n)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// adjust from 1-based indexing
k -= 1;
// Function Body
if (n <= 1)
return 0;
// zero out current top block of V on all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
magmablas_zlaset( MagmaUpperLower, nb, nb, dV(d,k,0), ldv );
}
// set all Y=0
lapackf77_zlaset( "Full", &n, &nb, &c_zero, &c_zero, Y, &ldy );
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Finish applying I - V * T * V' on right
tmp = MAGMA_Z_NEGATE( tau[i-1] );
blasf77_zaxpy( &n_k, &tmp, Y(k,i-1), &ione, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_zcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_ztrmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_zgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_ztrmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_zgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_ztrmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_zaxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_zlarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// compute yi = A vi = sum_g A{d} vi{d}
nblocks = (n-1) / nb / ngpu + 1;
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// dV(k+i+1:n-1, i) = VA(k+i:n, i)
magma_zsetvector_async( n_k_i_1,
A(k+i+1,i), 1,
dV(d, k+i+1, i), 1, data->streams[d] );
// copy column of dV -> dVd, using block cyclic distribution.
// This assumes V and Vd have been padded so that
// a 2D matrix copy doesn't access them out-of-bounds
gblock = k / nb;
lblock = gblock / ngpu;
lgid = gblock % ngpu;
if ( d < lgid ) {
lblock += 1;
}
// treat V as (nb*ngpu) x nblock matrix, and Vd as nb x nblock matrix
magmablas_zlacpy( 'F', nb, nblocks-lblock,
dV (d, d*nb + lblock*nb*ngpu, i), nb*ngpu,
dVd(d, 0 + lblock*nb , i), nb );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
// dY(k:n, i) = dA(k:n, k+i+1:n) * dV(k+i+1:n, i)
// skip if matrix is empty
// each GPU copies to different temporary vector in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_zgemv( 'N', n-k, dn-dki1,
c_one, dA (d, k , dki1), ldda,
dVd(d, dki1, i), 1,
c_zero, dY (d, k , i), 1 );
// copy vector to host, storing in column nb+d of Y
// as temporary space (Y has >= nb+ngpu columns)
magma_zgetvector_async( n-k,
dY(d, k, i), 1,
Y(k, nb+d), 1, data->streams[d] );
}
}
// while GPU is doing above Ag*v...
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_Z_NEGATE( tau[i] );
blasf77_zgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_ztrmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// apply reflectors to next column, A(i+1), on right only.
// one axpy will be required to finish this, in the next iteration above
if ( i > 0 && i+1 < nb ) {
// Update next column, A(k:n,i+1), applying Q on right.
// One axpy will be required to finish this, in the next iteration
// above, after yi is computed.
// This updates one more row than LAPACK does (row k),
// making block above panel an even multiple of nb.
// Use last column of T as workspace, w.
magma_int_t i1 = i+1;
// If complex, conjugate row of V, and undo afterwards
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i1, A(k+i1,0), &lda );
#endif
// w = T(0:i, 0:i+1) * VA(k+i+1, 0:i+1)'
// T is now rectangular, so we use gemv instead of trmv as in lapack.
blasf77_zgemv( "No trans", &i, &i1,
&c_one, T(0,0), &ldt,
A(k+i1,0), &lda,
&c_zero, T(0,nb-1), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_zgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i1), &ione );
}
// yi = sum_g yi{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( data->streams[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// yi = yi + yi{d}
blasf77_zaxpy( &n_k, &c_one, Y(k,nb+d), &ione, Y(k,i), &ione );
}
}
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
// compute Y = Am V = sum_g Am{d} V{d} --- top part, Y(0:k-1,:)
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
// dY(0:k, :) = dA(0:k, k+i+1:n-1) * dV(k+i+1:n-1, :)
// skip if matrix is empty
// each GPU copies to different temporary block in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_zgemm( 'N', 'N', k, nb, dn-dki1,
c_one, dA (d, 0 , dki1), ldda,
dVd(d, dki1, 0), ldvd,
c_zero, dY (d, 0 , 0), ldda );
// copy result to host, storing in columns [nb + nb*d : nb + nb*(d+1)] of Y
// as temporary space (Y has nb + nb*ngpu columns)
magma_zgetmatrix_async( k, nb,
dY(d, 0, 0), ldda,
Y(0,nb+nb*d), ldy, data->streams[d] );
}
}
// Y = sum_g Y{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( 0 );
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// Y = Y + Am V
for( i = 0; i < nb; ++i ) {
blasf77_zaxpy( &k, &c_one, Y(0,nb+nb*d+i), &ione, Y(0,i), &ione );
}
}
}
// copy Y and T matrices to GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_zsetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->streams[d] );
magma_zsetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->streams[d] );
}
return 0;
} // magma_zlahr2
extern "C" magma_int_t
magma_zlatrd2(char uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *a, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *w, magma_int_t ldw,
magmaDoubleComplex *da, magma_int_t ldda,
magmaDoubleComplex *dw, magma_int_t lddw,
magmaDoubleComplex *dwork, magma_int_t ldwork)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZLATRD2 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 = 'U', ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = 'L', ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the matrix A.
NB (input) INTEGER
The number of rows and columns to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', 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 = 'L', 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 = 'U', 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 = 'L', 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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= (1,N).
E (output) COMPLEX_16 array, dimension (N-1)
If UPLO = 'U', E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = 'L', E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
TAU (output) COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = 'U', and in TAU(1:nb) if UPLO = 'L'.
See Further Details.
W (output) COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
LDW (input) INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
===============
If UPLO = 'U', 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 = 'L', 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 = 'U': if UPLO = 'L':
( 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).
===================================================================== */
char uplo_[2] = {uplo, 0};
magma_int_t i;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_zmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside zlatrd
if (lapackf77_lsame(uplo_, "U")) {
/* 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_zlacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_zgemv("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_zlacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_zgemv("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_zlacgv(&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_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
MAGMA_Z_SET2REAL(*A(i-1, i), 1.);
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1 );
#if (GPUSHMEM < 200)
magma_zhemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione);
#else
magmablas_zhemv2(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_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_zgemv(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_zgemv("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_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("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_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy(&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_zlacgv(&i, W(i, 0), &ldw);
#endif
blasf77_zgemv("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_zlacgv(&i, W(i, 0), &ldw);
lapackf77_zlacgv(&i, A(i ,0), &lda);
#endif
blasf77_zgemv("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_zlacgv(&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_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_Z_REAL( alpha );
MAGMA_Z_SET2REAL(*A(i+1, i), 1.);
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
#if (GPUSHMEM < 200)
magma_zhemv(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_zhemv2('L', 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_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(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_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("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_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* zlatrd */
extern "C" magma_int_t
magma_zlahr2(magma_int_t n, magma_int_t k, magma_int_t nb,
cuDoubleComplex *da, cuDoubleComplex *dv,
cuDoubleComplex *a, magma_int_t lda,
cuDoubleComplex *tau, cuDoubleComplex *t, magma_int_t ldt,
cuDoubleComplex *y, magma_int_t ldy)
{
/* -- MAGMA auxiliary routine (version 1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
ZLAHR2 reduces the first NB columns of a complex general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
This is an auxiliary routine called by ZGEHRD.
Arguments
=========
N (input) INTEGER
The order of the matrix A.
K (input) INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
NB (input) INTEGER
The number of columns to be reduced.
DA (input/output) COMPLEX_16 array on the GPU, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
DV (output) COMPLEX_16 array on the GPU, dimension (N, NB)
On exit this contains the Householder vectors of the transformation.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX_16 array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) COMPLEX_16 array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) COMPLEX_16 array, dimension (LDY,NB)
The n-by-nb matrix Y.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
Further Details
===============
The matrix Q is represented as a product of nb 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+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
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.
===================================================================== */
cuDoubleComplex c_zero = MAGMA_Z_ZERO;
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ldda = lda;
magma_int_t c__1 = 1;
magma_int_t a_dim1, a_offset, t_dim1, t_offset, y_dim1, y_offset, i__2, i__3;
cuDoubleComplex d__1;
magma_int_t i__;
cuDoubleComplex ei;
--tau;
a_dim1 = lda;
a_offset = 1 + a_dim1;
a -= a_offset;
t_dim1 = ldt;
t_offset = 1 + t_dim1;
t -= t_offset;
y_dim1 = ldy;
y_offset = 1 + y_dim1;
y -= y_offset;
/* Function Body */
if (n <= 1)
return 0;
for (i__ = 1; i__ <= nb; ++i__) {
if (i__ > 1) {
/* Update A(K+1:N,I); Update I-th column of A - Y * V' */
i__2 = n - k + 1;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
blasf77_zcopy(&i__3, &a[k+i__-1+a_dim1], &lda, &t[nb*t_dim1+1], &c__1);
blasf77_ztrmv("u","n","n",&i__3,&t[t_offset], &ldt, &t[nb*t_dim1+1], &c__1);
blasf77_zgemv("NO TRANSPOSE", &i__2, &i__3, &c_neg_one, &y[k + y_dim1],
&ldy, &t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__*a_dim1],&c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
/* Apply I - V * T' * V' to this column (call it b) from the
left, using the last column of T as workspace
Let V = ( V1 ) and b = ( b1 ) (first I-1 rows)
( V2 ) ( b2 )
where V1 is unit lower triangular
w := V1' * b1 */
i__2 = i__ - 1;
blasf77_zcopy(&i__2, &a[k+1+i__*a_dim1], &c__1, &t[nb*t_dim1+1], &c__1);
blasf77_ztrmv("Lower", MagmaConjTransStr, "UNIT", &i__2,
&a[k + 1 + a_dim1], &lda, &t[nb * t_dim1 + 1], &c__1);
/* w := w + V2'*b2 */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_one, &t[nb*t_dim1+1], &c__1);
/* w := T'*w */
i__2 = i__ - 1;
blasf77_ztrmv("U", MagmaConjTransStr, "N", &i__2, &t[t_offset], &ldt,
&t[nb*t_dim1+1], &c__1);
/* b2 := b2 - V2*w */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv("N", &i__2, &i__3, &c_neg_one, &a[k + i__ + a_dim1], &lda,
&t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__+i__*a_dim1], &c__1);
/* b1 := b1 - V1*w */
i__2 = i__ - 1;
blasf77_ztrmv("L","N","U",&i__2,&a[k+1+a_dim1],&lda,&t[nb*t_dim1+1],&c__1);
blasf77_zaxpy(&i__2, &c_neg_one, &t[nb * t_dim1 + 1], &c__1,
&a[k + 1 + i__ * a_dim1], &c__1);
a[k + i__ - 1 + (i__ - 1) * a_dim1] = ei;
}
/* Generate the elementary reflector H(I) to annihilate A(K+I+1:N,I) */
i__2 = n - k - i__ + 1;
i__3 = k + i__ + 1;
lapackf77_zlarfg(&i__2, &a[k + i__ + i__ * a_dim1],
&a[min(i__3,n) + i__ * a_dim1], &c__1, &tau[i__]);
ei = a[k + i__ + i__ * a_dim1];
a[k + i__ + i__ * a_dim1] = c_one;
/* Compute Y(K+1:N,I) */
i__2 = n - k;
i__3 = n - k - i__ + 1;
magma_zsetvector( i__3,
&a[k + i__ + i__*a_dim1], 1,
dv+(i__-1)*(ldda+1), 1 );
magma_zgemv(MagmaNoTrans, i__2+1, i__3, c_one,
da -1 + k + i__ * ldda, ldda,
dv+(i__-1)*(ldda+1), c__1, c_zero,
da-1 + k + (i__-1)*ldda, c__1);
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_zgemv(MagmaConjTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_zero, &t[i__*t_dim1+1], &c__1);
/* Compute T(1:I,I) */
i__2 = i__ - 1;
d__1 = MAGMA_Z_NEGATE( tau[i__] );
blasf77_zscal(&i__2, &d__1, &t[i__ * t_dim1 + 1], &c__1);
blasf77_ztrmv("U","N","N", &i__2, &t[t_offset], &ldt, &t[i__*t_dim1+1], &c__1);
t[i__ + i__ * t_dim1] = tau[i__];
magma_zgetvector( n - k + 1,
da-1+ k+(i__-1)*ldda, 1,
y+ k + i__*y_dim1, 1 );
}
a[k + nb + nb * a_dim1] = ei;
return 0;
} /* magma_zlahr2 */
/**
Purpose
-------
ZLATRD 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, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@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]
nb0 INTEGER
The block size used for the matrix distribution.
nb and nb0 can be different for the final step of zhetrd.
@param[in,out]
A COMPLEX_16 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_16 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_16 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_16 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).
@param
dA
@param[in]
ldda
@param[in]
offset
@param
dW
@param[in]
lddw
@param
hwork
@param[in]
lhwork
@param
dwork
@param[in]
ldwork
@param[in]
queues magma_queue_t array of dimension (ngpu).
queues[dev] is an execution queue on GPU dev.
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_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo,
magma_int_t n, magma_int_t nb, magma_int_t nb0,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex_ptr dA[], magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex_ptr dW[], magma_int_t lddw,
magmaDoubleComplex *hwork, magma_int_t lhwork,
magmaDoubleComplex_ptr dwork[], magma_int_t ldwork,
magma_queue_t queues[] )
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(dev, i, j) (dA[(dev)] + ((j)+loffset)*ldda + (i) + offset)
#define dW(dev, i, j) (dW[(dev)] + (j) *lddw + (i))
#define dW1(dev, i, j) (dW[(dev)] + ((j)+nb) *lddw + (i))
/* Constants */
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magma_int_t ione = 1;
/* Local variables */
magmaDoubleComplex alpha, value;
magma_int_t dev;
magma_int_t i, n_i, n_i_1, ip1, iw;
// TODO check arguments
magma_int_t info = 0;
if (n <= 0) {
return info;
}
// TODO allocate f in zhetrd and pass into zlatrd. (e.g., expand hwork a bit)
magmaDoubleComplex *f;
magma_zmalloc_cpu( &f, n );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
return info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
ip1 = i + 1;
n_i_1 = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
magmaDoubleComplex wii = -conj( *W(i, iw+1) );
blasf77_zaxpy( &ip1, &wii, A(0, i+1), &ione, A(0, i), &ione );
wii = -conj( *A(i, i+1) );
blasf77_zaxpy( &ip1, &wii, W(0, iw+1), &ione, A(0, i), &ione );
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_ONE;
// TODO Previously, this set dx2[dev] = dW1(dev, 0, iw); and used dx2 in zhemv.
// TODO Now zhemv handles broadcasting x to the GPUs, but data in dW1 is
// TODO apparently still used in zhetrd_mgpu / zher2k_mgpu.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, A(0,i), 1, dW1(dev, 0, iw), 1, queues[dev] );
}
magmablas_zhemv_mgpu(
MagmaUpper, i, c_one, dA, ldda, 0,
A(0,i), 1, c_zero, W(0, iw), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
if (i < n-1) {
blasf77_zgemv( MagmaConjTransStr, &i, &n_i_1, &c_one,
W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero,
W(i+1, iw), &ione );
}
/* overlap update */
if ( i < n-1 && i-1 >= n - nb ) {
/* Update A(1:i,i) */
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, W(i-1, iw+1), &ldw );
#endif
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
A(0, i+1), &lda,
W(i-1, iw+1), &ldw, &c_one,
A(0, i-1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, W(i-1, iw+1), &ldw );
lapackf77_zlacgv( &n_i_1, A(i-1, i +1), &lda );
#endif
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
W(0, iw+1), &ldw,
A(i-1, i+1), &lda, &c_one,
A(0, i-1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, A(i-1, i+1), &lda );
#endif
}
// synchronize to get zhemv result W(0, iw)
magmablas_zhemv_mgpu_sync(
MagmaUpper, i, c_one, dA, ldda, 0,
A(0,i), 1, c_zero, W(0, iw), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
if (i < n-1) {
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one,
W(0, iw), &ione );
blasf77_zgemv( MagmaConjTransStr, &i, &n_i_1, &c_one,
A(0, i+1), &lda,
A(0, i), &ione, &c_zero,
W(i+1, iw), &ione );
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one,
W(0, iw), &ione );
}
blasf77_zscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy( &i, &alpha, A(0, i), &ione, W(0, iw), &ione );
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, W(0,iw), 1, dW(dev, 0, iw), 1, queues[dev] );
}
}
}
} else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
n_i = n - i;
//idw = ((offset+i)/nb)%ngpu;
if ( i > 0 ) {
trace_cpu_start( 0, "gemv", "gemv" );
magmaDoubleComplex wii = -conj( *W(i, i-1) );
blasf77_zaxpy( &n_i, &wii, A(i, i-1), &ione, A(i, i), &ione );
wii = -conj( *A(i, i-1) );
blasf77_zaxpy( &n_i, &wii, W(i, i-1), &ione, A(i, i), &ione );
}
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
n_i_1 = n - i - 1;
trace_cpu_start( 0, "larfg", "larfg" );
alpha = *A(i+1, i);
lapackf77_zlarfg( &n_i_1, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_ONE;
trace_cpu_end( 0 );
/* Compute W(i+1:n,i) */
// TODO Previously, this set dx2[id] = dW1(id, 0, i)-offset; and used dx2 in zhemv.
// TODO Now zhemv handles broadcasting x to the GPUs, but data in dW1 is
// TODO apparently still used in zhetrd_mgpu / zher2k_mgpu.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, A(0,i), 1, dW1(dev, 0, i), 1, queues[dev] );
}
magmablas_zhemv_mgpu(
MagmaLower, n_i_1, c_one, dA, ldda, offset+i+1,
A(i+1, i), 1, c_zero, W(i+1, i), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
trace_cpu_start( 0, "gemv", "gemv" );
blasf77_zgemv( MagmaConjTransStr, &n_i_1, &i, &c_one,
W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero,
W(0, i), &ione );
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
A(i+1, 0), &lda,
W(0, i), &ione, &c_zero,
f, &ione );
blasf77_zgemv( MagmaConjTransStr, &n_i_1, &i, &c_one,
A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero,
W(0, i), &ione );
trace_cpu_end( 0 );
/* overlap update */
if ( i > 0 && i+1 < n ) {
trace_cpu_start( 0, "gemv", "gemv" );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i+1, 0), &ldw );
#endif
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
A(i+1, 0), &lda,
W(i+1, 0), &ldw, &c_one,
A(i+1, i+1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i+1, 0), &ldw );
lapackf77_zlacgv( &i, A(i+1, 0), &lda );
#endif
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
W(i+1, 0), &ldw,
A(i+1, 0), &lda, &c_one,
A(i+1, i+1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, A(i+1, 0), &lda );
#endif
trace_cpu_end( 0 );
}
// synchronize to get zhemv result W(i+1, i)
magmablas_zhemv_mgpu_sync(
MagmaLower, n_i_1, c_one, dA, ldda, offset+i+1,
A(i+1, i), 1, c_zero, W(i+1, i), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
trace_cpu_start( 0, "axpy", "axpy" );
if (i != 0) {
blasf77_zaxpy( &n_i_1, &c_one, f, &ione, W(i+1, i), &ione );
}
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
W(i+1, 0), &ldw,
W(0, i), &ione, &c_one,
W(i+1, i), &ione );
blasf77_zscal( &n_i_1, &tau[i], W(i+1,i), &ione );
value = magma_cblas_zdotc( n_i_1, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy( &n_i_1, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
trace_cpu_end( 0 );
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, W(0,i), 1, dW(dev, 0, i), 1, queues[dev] );
}
}
}
}
magma_free_cpu( f );
magma_setdevice( orig_dev );
return info;
} /* magma_zlatrd_mgpu */
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
double magma_error, cublas_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ycublas, *Ymagma;
magmaDoubleComplex *dA, *dX, *dY;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("trans = %s\n", lapack_trans_const(opts.transA) );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = ((M+31)/32)*32;
gflops = FLOPS_ZGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ycublas, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, sizeA );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, sizeY );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &sizeA, A );
lapackf77_zlarnv( &ione, ISEED, &sizeX, X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( M, N, A, lda, dA, lda );
magma_zsetvector( Xm, X, incx, dX, incx );
magma_zsetvector( Ym, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasZgemv( handle, cublas_trans_const(opts.transA),
M, N, &alpha, dA, lda, dX, incx, &beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetvector( Ym, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_zsetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_zgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( Ym, dY, incy, Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_zgemv( lapack_trans_const(opts.transA), &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_zaxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_zlange( "M", &Ym, &ione, Ymagma, &Ym, work ) / Ym;
blasf77_zaxpy( &Ym, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_zlange( "M", &Ym, &ione, Ycublas, &Ym, work ) / Ym;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error,
(magma_error < tol && cublas_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && cublas_error < tol);
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgemm_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cpu_perf, cpu_time;
double magma_error, magma_err, Ynorm, work[1];
magma_int_t M, N, Xm, Ym, lda, ldda;
magma_int_t sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_int_t batchCount;
magmaDoubleComplex *h_A, *h_X, *h_Y, *h_Ymagma;
magmaDoubleComplex *d_A, *d_X, *d_Y;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.48, 0.38 );
magmaDoubleComplex **A_array = NULL;
magmaDoubleComplex **X_array = NULL;
magmaDoubleComplex **Y_array = NULL;
magma_opts opts;
parse_opts( argc, argv, &opts );
batchCount = opts.batchcount;
opts.lapack |= opts.check;
//double tol = opts.tolerance * lapackf77_dlamch("E");
printf("trans = %s\n", lapack_trans_const(opts.transA) );
printf("BatchCount M N MAGMA Gflop/s (ms) CPU Gflop/s (ms) MAGMA error\n");
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = ((M+31)/32)*32;
gflops = FLOPS_ZGEMV( M, N ) / 1e9 * batchCount;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N*batchCount;
sizeX = incx*Xm*batchCount;
sizeY = incy*Ym*batchCount;
ldda = ((lda+31)/32)*32;
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( h_X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( h_Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( h_Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N*batchCount );
TESTING_MALLOC_DEV( d_X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( d_Y, magmaDoubleComplex, sizeY );
magma_malloc((void**)&A_array, batchCount * sizeof(*A_array));
magma_malloc((void**)&X_array, batchCount * sizeof(*X_array));
magma_malloc((void**)&Y_array, batchCount * sizeof(*Y_array));
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_zlarnv( &ione, ISEED, &sizeX, h_X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, h_Y );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_zsetmatrix( M, N*batchCount, h_A, lda, d_A, ldda );
magma_zsetvector( Xm*batchCount, h_X, incx, d_X, incx );
magma_zsetvector( Ym*batchCount, h_Y, incy, d_Y, incy );
zset_pointer(A_array, d_A, ldda, 0, 0, ldda*N, batchCount, magma_stream);
zset_pointer(X_array, d_X, 1, 0, 0, incx*Xm, batchCount, magma_stream);
zset_pointer(Y_array, d_Y, 1, 0, 0, incy*Ym, batchCount, magma_stream);
magma_time = magma_sync_wtime( NULL );
magmablas_zgemv_batched(opts.transA, M, N,
alpha, A_array, ldda,
X_array, incx,
beta, Y_array, incy, batchCount, magma_stream);
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( Ym*batchCount, d_Y, incy, h_Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
for(int i=0; i<batchCount; i++)
{
blasf77_zgemv(
lapack_trans_const(opts.transA),
&M, &N,
&alpha, h_A + i*lda*N, &lda,
h_X + i*Xm, &incx,
&beta, h_Y + i*Ym, &incy );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
magma_error = 0.0;
for(int s=0; s<batchCount; s++)
{
Ynorm = lapackf77_zlange( "M", &M, &ione, h_Y + s*Ym, &incy, work );
blasf77_zaxpy( &Ym, &c_neg_one, h_Y + s*Ym, &ione, h_Ymagma + s*Ym, &ione );
magma_err = lapackf77_zlange( "M", &M, &ione, h_Ymagma + s*Ym, &incy, work ) / Ynorm;
if ( isnan(magma_err) || isinf(magma_err) ) {
magma_error = magma_err;
break;
}
magma_error = max(fabs(magma_err), magma_error);
}
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e \n",
(int) batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time,
cpu_perf, 1000.*cpu_time,
magma_error);
}
else {
printf("%10d %5d %5d %7.2f (%7.2f) --- ( --- ) ---\n",
(int) batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time);
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_Y );
TESTING_FREE_CPU( h_Ymagma );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_X );
TESTING_FREE_DEV( d_Y );
TESTING_FREE_DEV( A_array );
TESTING_FREE_DEV( X_array );
TESTING_FREE_DEV( Y_array );
fflush( stdout);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
ZLATRD2 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, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD2_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_16 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_16 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_16 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_16 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_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd2(magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dW, magma_int_t lddw,
magmaDoubleComplex *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;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_zmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside zlatrd
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_zlacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_zgemv("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_zlacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_zgemv("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_zlacgv(&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_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_MAKE( 1, 0 );
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1 );
//#if (GPUSHMEM < 200)
//magma_zhemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
// dA(0, i), ione, c_zero, dW(0, iw), ione);
//#else
magmablas_zhemv_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_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_zgemv(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_zgemv("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_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("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_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy(&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_zlacgv(&i, W(i, 0), &ldw);
#endif
blasf77_zgemv("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_zlacgv(&i, W(i, 0), &ldw);
lapackf77_zlacgv(&i, A(i, 0), &lda);
#endif
blasf77_zgemv("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_zlacgv(&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_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_MAKE( 1, 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
//#if (GPUSHMEM < 200)
//magma_zhemv(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_zhemv_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_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(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_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("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_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* magma_zlatrd */
magma_int_t magma_ztrevc3_mt(
magma_side_t side, magma_vec_t howmany,
magma_int_t *select, // logical in Fortran
magma_int_t n,
magmaDoubleComplex *T, magma_int_t ldt,
magmaDoubleComplex *VL, magma_int_t ldvl,
magmaDoubleComplex *VR, magma_int_t ldvr,
magma_int_t mm, magma_int_t *mout,
magmaDoubleComplex *work, magma_int_t lwork,
#ifdef COMPLEX
double *rwork,
#endif
magma_int_t *info )
{
#define T(i,j) ( T + (i) + (j)*ldt )
#define VL(i,j) (VL + (i) + (j)*ldvl)
#define VR(i,j) (VR + (i) + (j)*ldvr)
#define work(i,j) (work + (i) + (j)*n)
// .. Parameters ..
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magma_int_t nbmin = 16, nbmax = 128;
const magma_int_t ione = 1;
// .. Local Scalars ..
magma_int_t allv, bothv, leftv, over, rightv, somev;
magma_int_t i, ii, is, j, k, ki, iv, n2, nb, nb2, version;
double ovfl, remax, unfl; //smlnum, smin, ulp
// Decode and test the input parameters
bothv = (side == MagmaBothSides);
rightv = (side == MagmaRight) || bothv;
leftv = (side == MagmaLeft ) || bothv;
allv = (howmany == MagmaAllVec);
over = (howmany == MagmaBacktransVec);
somev = (howmany == MagmaSomeVec);
// Set mout to the number of columns required to store the selected
// eigenvectors.
if ( somev ) {
*mout = 0;
for( j=0; j < n; ++j ) {
if ( select[j] ) {
*mout += 1;
}
}
}
else {
*mout = n;
}
*info = 0;
if ( ! rightv && ! leftv )
*info = -1;
else if ( ! allv && ! over && ! somev )
*info = -2;
else if ( n < 0 )
*info = -4;
else if ( ldt < max( 1, n ) )
*info = -6;
else if ( ldvl < 1 || ( leftv && ldvl < n ) )
*info = -8;
else if ( ldvr < 1 || ( rightv && ldvr < n ) )
*info = -10;
else if ( mm < *mout )
*info = -11;
else if ( lwork < max( 1, 2*n ) )
*info = -14;
if ( *info != 0 ) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// Quick return if possible.
if ( n == 0 ) {
return *info;
}
// Use blocked version (2) if sufficient workspace.
// Requires 1 vector to save diagonal elements, and 2*nb vectors for x and Q*x.
// (Compared to dtrevc3, rwork stores 1-norms.)
// Zero-out the workspace to avoid potential NaN propagation.
nb = 2;
if ( lwork >= n + 2*n*nbmin ) {
version = 2;
nb = (lwork - n) / (2*n);
nb = min( nb, nbmax );
nb2 = 1 + 2*nb;
lapackf77_zlaset( "F", &n, &nb2, &c_zero, &c_zero, work, &n );
}
else {
version = 1;
}
// Set the constants to control overflow.
unfl = lapackf77_dlamch( "Safe minimum" );
ovfl = 1. / unfl;
lapackf77_dlabad( &unfl, &ovfl );
//ulp = lapackf77_dlamch( "Precision" );
//smlnum = unfl*( n / ulp );
// Store the diagonal elements of T in working array work.
for( i=0; i < n; ++i ) {
*work(i,0) = *T(i,i);
}
// Compute 1-norm of each column of strictly upper triangular
// part of T to control overflow in triangular solver.
rwork[0] = 0.;
for( j=1; j < n; ++j ) {
rwork[j] = magma_cblas_dzasum( j, T(0,j), ione );
}
// launch threads -- each single-threaded MKL
magma_int_t nthread = magma_get_parallel_numthreads();
magma_int_t lapack_nthread = magma_get_lapack_numthreads();
magma_set_lapack_numthreads( 1 );
magma_thread_queue queue;
queue.launch( nthread );
//printf( "nthread %d, %d\n", nthread, lapack_nthread );
// gemm_nb = N/thread, rounded up to multiple of 16,
// but avoid multiples of page size, e.g., 512*8 bytes = 4096.
magma_int_t gemm_nb = magma_int_t( ceil( ceil( ((double)n) / nthread ) / 16. ) * 16. );
if ( gemm_nb % 512 == 0 ) {
gemm_nb += 32;
}
magma_timer_t time_total=0, time_trsv=0, time_gemm=0, time_gemv=0, time_trsv_sum=0, time_gemm_sum=0, time_gemv_sum=0;
timer_start( time_total );
if ( rightv ) {
// ============================================================
// Compute right eigenvectors.
// iv is index of column in current block.
// Non-blocked version always uses iv=1;
// blocked version starts with iv=nb, goes down to 1.
// (Note the "0-th" column is used to store the original diagonal.)
iv = 1;
if ( version == 2 ) {
iv = nb;
}
timer_start( time_trsv );
is = *mout - 1;
for( ki=n-1; ki >= 0; --ki ) {
if ( somev ) {
if ( ! select[ki] ) {
continue;
}
}
//smin = max( ulp*MAGMA_Z_ABS1( *T(ki,ki) ), smlnum );
// --------------------------------------------------------
// Complex right eigenvector
*work(ki,iv) = c_one;
// Form right-hand side.
for( k=0; k < ki; ++k ) {
*work(k,iv) = -(*T(k,ki));
}
// Solve upper triangular system:
// [ T(1:ki-1,1:ki-1) - T(ki,ki) ]*X = scale*work.
if ( ki > 0 ) {
queue.push_task( new magma_zlatrsd_task(
MagmaUpper, MagmaNoTrans, MagmaNonUnit, MagmaTrue,
ki, T, ldt, *T(ki,ki),
work(0,iv), work(ki,iv), rwork ));
}
// Copy the vector x or Q*x to VR and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VR and normalize
queue.sync();
n2 = ki+1;
blasf77_zcopy( &n2, work(0,iv), &ione, VR(0,is), &ione );
ii = blasf77_izamax( &n2, VR(0,is), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *VR(ii,is) );
blasf77_zdscal( &n2, &remax, VR(0,is), &ione );
for( k=ki+1; k < n; ++k ) {
*VR(k,is) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
queue.sync();
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemv );
if ( ki > 0 ) {
blasf77_zgemv( "n", &n, &ki, &c_one,
VR, &ldvr,
work(0, iv), &ione,
work(ki,iv), VR(0,ki), &ione );
}
time_gemv_sum += timer_stop( time_gemv );
ii = blasf77_izamax( &n, VR(0,ki), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *VR(ii,ki) );
blasf77_zdscal( &n, &remax, VR(0,ki), &ione );
timer_start( time_trsv );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out below vector
for( k=ki+1; k < n; ++k ) {
*work(k,iv) = c_zero;
}
// Columns iv:nb of work are valid vectors.
// When the number of vectors stored reaches nb,
// or if this was last vector, do the GEMM
if ( (iv == 1) || (ki == 0) ) {
queue.sync();
time_trsv_sum += timer_stop( time_trsv );
timer_start( time_gemm );
nb2 = nb-iv+1;
n2 = ki+nb-iv+1;
// split gemm into multiple tasks, each doing one block row
for( i=0; i < n; i += gemm_nb ) {
magma_int_t ib = min( gemm_nb, n-i );
queue.push_task( new zgemm_task(
MagmaNoTrans, MagmaNoTrans, ib, nb2, n2, c_one,
VR(i,0), ldvr,
work(0,iv ), n, c_zero,
work(i,nb+iv), n ));
}
queue.sync();
time_gemm_sum += timer_stop( time_gemm );
// normalize vectors
// TODO if somev, should copy vectors individually to correct location.
for( k = iv; k <= nb; ++k ) {
ii = blasf77_izamax( &n, work(0,nb+k), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *work(ii,nb+k) );
blasf77_zdscal( &n, &remax, work(0,nb+k), &ione );
}
lapackf77_zlacpy( "F", &n, &nb2, work(0,nb+iv), &n, VR(0,ki), &ldvr );
iv = nb;
timer_start( time_trsv );
}
else {
iv -= 1;
}
} // blocked back-transform
is -= 1;
}
}
timer_stop( time_trsv );
timer_stop( time_total );
timer_printf( "trevc trsv %.4f, gemm %.4f, gemv %.4f, total %.4f\n",
time_trsv_sum, time_gemm_sum, time_gemv_sum, time_total );
if ( leftv ) {
// ============================================================
// Compute left eigenvectors.
// iv is index of column in current block.
// Non-blocked version always uses iv=1;
// blocked version starts with iv=1, goes up to nb.
// (Note the "0-th" column is used to store the original diagonal.)
iv = 1;
is = 0;
for( ki=0; ki < n; ++ki ) {
if ( somev ) {
if ( ! select[ki] ) {
continue;
}
}
//smin = max( ulp*MAGMA_Z_ABS1( *T(ki,ki) ), smlnum );
// --------------------------------------------------------
// Complex left eigenvector
*work(ki,iv) = c_one;
// Form right-hand side.
for( k = ki + 1; k < n; ++k ) {
*work(k,iv) = -MAGMA_Z_CONJ( *T(ki,k) );
}
// Solve conjugate-transposed triangular system:
// [ T(ki+1:n,ki+1:n) - T(ki,ki) ]**H * X = scale*work.
// TODO what happens with T(k,k) - lambda is small? Used to have < smin test.
if ( ki < n-1 ) {
n2 = n-ki-1;
queue.push_task( new magma_zlatrsd_task(
MagmaUpper, MagmaConjTrans, MagmaNonUnit, MagmaTrue,
n2, T(ki+1,ki+1), ldt, *T(ki,ki),
work(ki+1,iv), work(ki,iv), rwork ));
}
// Copy the vector x or Q*x to VL and normalize.
if ( ! over ) {
// ------------------------------
// no back-transform: copy x to VL and normalize
queue.sync();
n2 = n-ki;
blasf77_zcopy( &n2, work(ki,iv), &ione, VL(ki,is), &ione );
ii = blasf77_izamax( &n2, VL(ki,is), &ione ) + ki - 1;
remax = 1. / MAGMA_Z_ABS1( *VL(ii,is) );
blasf77_zdscal( &n2, &remax, VL(ki,is), &ione );
for( k=0; k < ki; ++k ) {
*VL(k,is) = c_zero;
}
}
else if ( version == 1 ) {
// ------------------------------
// version 1: back-transform each vector with GEMV, Q*x.
queue.sync();
if ( ki < n-1 ) {
n2 = n-ki-1;
blasf77_zgemv( "n", &n, &n2, &c_one,
VL(0,ki+1), &ldvl,
work(ki+1,iv), &ione,
work(ki, iv), VL(0,ki), &ione );
}
ii = blasf77_izamax( &n, VL(0,ki), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *VL(ii,ki) );
blasf77_zdscal( &n, &remax, VL(0,ki), &ione );
}
else if ( version == 2 ) {
// ------------------------------
// version 2: back-transform block of vectors with GEMM
// zero out above vector
// could go from (ki+1)-NV+1 to ki
for( k=0; k < ki; ++k ) {
*work(k,iv) = c_zero;
}
// Columns 1:iv of work are valid vectors.
// When the number of vectors stored reaches nb,
// or if this was last vector, do the GEMM
if ( (iv == nb) || (ki == n-1) ) {
queue.sync();
n2 = n-(ki+1)+iv;
// split gemm into multiple tasks, each doing one block row
for( i=0; i < n; i += gemm_nb ) {
magma_int_t ib = min( gemm_nb, n-i );
queue.push_task( new zgemm_task(
MagmaNoTrans, MagmaNoTrans, ib, iv, n2, c_one,
VL(i,ki-iv+1), ldvl,
work(ki-iv+1,1), n, c_zero,
work(i,nb+1), n ));
}
queue.sync();
// normalize vectors
for( k=1; k <= iv; ++k ) {
ii = blasf77_izamax( &n, work(0,nb+k), &ione ) - 1;
remax = 1. / MAGMA_Z_ABS1( *work(ii,nb+k) );
blasf77_zdscal( &n, &remax, work(0,nb+k), &ione );
}
lapackf77_zlacpy( "F", &n, &iv, work(0,nb+1), &n, VL(0,ki-iv+1), &ldvl );
iv = 1;
}
else {
iv += 1;
}
} // blocked back-transform
is += 1;
}
}
// close down threads
queue.quit();
magma_set_lapack_numthreads( lapack_nthread );
return *info;
} // End of ZTREVC
/**
Purpose
-------
ZLATRD 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, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
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_16 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_16 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_16 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_16 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).
@param
dA TODO: dimension (ldda, n)?
@param
ldda TODO: ldda >= n?
@param
dW TODO: dimension (lddw, ??)
@param
lddw TODO: lddw >= n ??
@param[in]
queue magma_queue_t
Queue to execute in.
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_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex *work, magma_int_t lwork,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_ptr dW, magma_int_t lddw,
magma_queue_t queue )
{
#define A(i_, j_) (A + (i_) + (j_)*lda)
#define W(i_, j_) (W + (i_) + (j_)*ldw)
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#define dW(i_, j_) (dW + (i_) + (j_)*lddw)
/* Constants */
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magma_int_t ione = 1;
/* Local variables */
magmaDoubleComplex alpha, value;
magma_int_t i, i_n, i_1, iw;
/* Check arguments */
magma_int_t info = 0;
if ( uplo != MagmaLower && uplo != MagmaUpper ) {
info = -1;
} else if ( n < 0 ) {
info = -2;
} else if ( nb < 1 ) {
info = -3;
} else if ( lda < max(1,n) ) {
info = -5;
} else if ( ldw < max(1,n) ) {
info = -9;
} else if ( ldda < max(1,n) ) {
info = -11;
} else if ( lddw < max(1,n) ) {
info = -13;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Quick return if possible */
if (n == 0) {
return info;
}
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) */
#ifdef COMPLEX
lapackf77_zlacgv( &i_n, W(i, iw+1), &ldw );
#endif
blasf77_zgemv( "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 );
#ifdef COMPLEX
lapackf77_zlacgv( &i_n, W(i, iw+1), &ldw );
lapackf77_zlacgv( &i_n, A(i, i+1), &lda );
#endif
blasf77_zgemv( "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 );
#ifdef COMPLEX
lapackf77_zlacgv( &i_n, A(i, i+1), &lda );
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_ONE;
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1, queue );
magma_zhemv( MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione, queue );
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw), ldw, queue );
if (i < n-1) {
blasf77_zgemv( 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( queue );
if (i < n-1) {
blasf77_zgemv( "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_zgemv( MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
blasf77_zgemv( "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_zscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy( &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;
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i, 0), &ldw );
#endif
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i, 0), &ldw );
lapackf77_zlacgv( &i, A(i, 0), &lda );
#endif
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &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_zlarfg( &i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_ONE;
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1, queue );
magma_zhemv( MagmaLower, i_n, c_one, dA(i+1, i+1), ldda,
dA(i+1, i), ione, c_zero, dW(i+1, i), ione, queue );
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, queue );
blasf77_zgemv( MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, work, &ione );
blasf77_zgemv( 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( queue );
if (i != 0)
blasf77_zaxpy( &i_n, &c_one, work, &ione, W(i+1, i), &ione );
blasf77_zgemv( "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_zscal( &i_n, &tau[i], W(i+1,i), &ione );
value = magma_cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy( &i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
}
}
}
return info;
} /* magma_zlatrd */
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
double magma_error, dev_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, ldda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ydev, *Ymagma;
magmaDoubleComplex_ptr dA, dX, dY;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% trans = %s\n", lapack_trans_const(opts.transA) );
#ifdef HAVE_CUBLAS
printf("%% M N MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
printf("%% M N %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("%%==================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = M;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = FLOPS_ZGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ydev, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, sizeY );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &sizeA, A );
lapackf77_zlarnv( &ione, ISEED, &sizeX, X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( M, N, A, lda, dA, ldda, opts.queue );
magma_zsetvector( Xm, X, incx, dX, incx, opts.queue );
magma_zsetvector( Ym, Y, incy, dY, incy, opts.queue );
dev_time = magma_sync_wtime( opts.queue );
#ifdef HAVE_CUBLAS
cublasZgemv( opts.handle, cublas_trans_const(opts.transA),
M, N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
#else
magma_zgemv( opts.transA, M, N,
alpha, dA, ldda,
dX, incx,
beta, dY, incy );
#endif
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
dev_perf = gflops / dev_time;
magma_zgetvector( Ym, dY, incy, Ydev, incy, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_zsetvector( Ym, Y, incy, dY, incy, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
magmablas_zgemv( opts.transA, M, N, alpha, dA, ldda, dX, incx, beta, dY, incy, opts.queue );
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( Ym, dY, incy, Ymagma, incy, opts.queue );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_zgemv( lapack_trans_const(opts.transA), &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
double Anorm = lapackf77_zlange( "F", &M, &N, A, &lda, work );
double Xnorm = lapackf77_zlange( "F", &Xm, &ione, X, &Xm, work );
blasf77_zaxpy( &Ym, &c_neg_one, Y, &incy, Ydev, &incy );
dev_error = lapackf77_zlange( "F", &Ym, &ione, Ydev, &Ym, work ) / (Anorm * Xnorm);
#ifdef HAVE_CUBLAS
blasf77_zaxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_zlange( "F", &Ym, &ione, Ymagma, &Ym, work ) / (Anorm * Xnorm);
bool okay = (magma_error < tol) && (dev_error < tol);
status += ! okay;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(okay ? "ok" : "failed"));
#else
bool okay = (dev_error < tol);
status += ! okay;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(okay ? "ok" : "failed"));
#endif
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ydev );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
extern "C" double
magma_zlatrd_mgpu(magma_int_t num_gpus, char uplo,
magma_int_t n0, magma_int_t n, magma_int_t nb, magma_int_t nb0,
magmaDoubleComplex *a, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *w, magma_int_t ldw,
magmaDoubleComplex **da, magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex **dw, magma_int_t lddw,
magmaDoubleComplex *dwork[MagmaMaxGPUs], magma_int_t ldwork,
magma_int_t k,
magmaDoubleComplex *dx[MagmaMaxGPUs], magmaDoubleComplex *dy[MagmaMaxGPUs],
magmaDoubleComplex *work,
magma_queue_t stream[][10],
double *times)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZLATRD 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 = 'U', ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = 'L', ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the matrix A.
NB (input) INTEGER
The number of rows and columns to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', 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 = 'L', 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 = 'U', 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 = 'L', 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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= (1,N).
E (output) COMPLEX_16 array, dimension (N-1)
If UPLO = 'U', E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = 'L', E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
TAU (output) COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = 'U', and in TAU(1:nb) if UPLO = 'L'.
See Further Details.
W (output) COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
LDW (input) INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
===============
If UPLO = 'U', 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 = 'L', 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 = 'U': if UPLO = 'L':
( 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).
===================================================================== */
char uplo_[2] = {uplo, 0};
double mv_time = 0.0;
magma_int_t i;
#ifndef MAGMABLAS_ZHEMV_MGPU
magma_int_t loffset = nb0*((offset/nb0)/num_gpus);
#endif
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t id, idw, i_one = 1;
//magma_int_t kk;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *dx2[MagmaMaxGPUs];
magmaDoubleComplex *f = (magmaDoubleComplex *)malloc(n*sizeof(magmaDoubleComplex ));
if (n <= 0) {
return 0;
}
//#define PROFILE_SYMV
#ifdef PROFILE_SYMV
magma_event_t start, stop;
float etime;
magma_timestr_t cpu_start, cpu_end;
magma_setdevice(0);
magma_event_create( &start );
magma_event_create( &stop );
#endif
if (lapackf77_lsame(uplo_, "U")) {
/* 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) */
magmaDoubleComplex wii = *W(i, iw+1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_zaxpy(&i_1, &wii, A(0, i+1), &i_one, A(0, i), &ione);
wii = *A(i, i+1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_zaxpy(&i_1, &wii, W(0, iw+1), &i_one, A(0, i), &ione);
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_MAKE( 1, 0 );
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
dx2[id] = dW1(id, 0, iw);
magma_zsetvector_async( n, A(0,i), 1, dW1(id, 0, iw), 1, stream[id][0]);
#ifndef MAGMABLAS_ZHEMV_MGPU
magma_zsetvector_async( i, A(0,i), 1, dx[id], 1, stream[id][0] );
#endif
}
magmablas_zhemv_mgpu(num_gpus, k, 'U', i, nb0, c_one, da, ldda, 0,
dx2, ione, c_zero, dy, ione, dwork, ldwork,
work, W(0, iw), stream );
if (i < n-1) {
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
/* overlap update */
if( i < n-1 && i-1 >= n - nb )
{
magma_int_t im1_1 = i_1 - 1;
magma_int_t im1 = i-1;
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magma_int_t im1_n = i_n + 1;
lapackf77_zlacgv(&im1_n, W(im1, iw+1), &ldw);
#endif
blasf77_zgemv("No transpose", &im1_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(im1, iw+1), &ldw, &c_one, A(0, i-1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&im1_n, W(im1, iw+1), &ldw);
lapackf77_zlacgv(&im1_n, A(im1, i +1), &lda);
#endif
blasf77_zgemv("No transpose", &im1_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(im1, i+1), &lda, &c_one, A(0, i-1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&im1_n, A(im1, i+1), &lda);
#endif
}
// 3. Here is where we need it // TODO find the right place
magmablas_zhemv_sync(num_gpus, k, i, work, W(0, iw), stream );
if (i < n-1) {
blasf77_zgemv("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_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("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_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -.5f * value;
blasf77_zaxpy(&i, &alpha, A(0, i), &ione, W(0, iw), &ione);
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
if( k > 1 ) {
magma_zsetvector_async( n, W(0,iw), 1, dW(id, 0, iw), 1, stream[id][1] );
} else {
magma_zsetvector_async( n, W(0,iw), 1, dW(id, 0, iw), 1, stream[id][0] );
}
}
}
}
} else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
idw = ((offset+i)/nb)%num_gpus;
if( i > 0 ) {
trace_cpu_start( 0, "gemv", "gemv" );
magmaDoubleComplex wii = *W(i, i-1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_zaxpy( &i_n, &wii, A(i, i-1), &ione, A(i, i), &ione);
wii = *A(i, i-1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &lda);
#endif
wii = -wii;
blasf77_zaxpy( &i_n, &wii, W(i, i-1), &ione, A(i, i), &ione);
}
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
trace_cpu_start( 0, "larfg", "larfg" );
alpha = *A(i+1, i);
#ifdef PROFILE_SYMV
cpu_start = get_current_time();
#endif
lapackf77_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
#ifdef PROFILE_SYMV
cpu_end = get_current_time();
times[0] += GetTimerValue(cpu_start,cpu_end)/1000.0;
#endif
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_MAKE( 1, 0 );
trace_cpu_end( 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
//trace_gpu_start( idw, 0, "comm", "comm1" );
#ifndef MAGMABLAS_ZHEMV_MGPU
magma_setdevice(idw);
magma_zsetvector( i_n, A(i+1,i), 1, dA(idw, i+1, i), 1 );
#endif
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
trace_gpu_start( id, 0, "comm", "comm" );
#ifdef MAGMABLAS_ZHEMV_MGPU
dx2[id] = dW1(id, 0, i)-offset;
#else
dx2[id] = dx[id];
magma_zsetvector( i_n, A(i+1,i), 1, dx[id], 1 );
#endif
magma_zsetvector_async( n, A(0,i), 1, dW1(id, 0, i), 1, stream[id][0] );
trace_gpu_end( id, 0 );
}
/* mat-vec on multiple GPUs */
#ifdef PROFILE_SYMV
magma_setdevice(0);
magma_event_record(start, stream[0][0]);
#endif
magmablas_zhemv_mgpu(num_gpus, k, 'L', i_n, nb0, c_one, da, ldda, offset+i+1,
dx2, ione, c_zero, dy, ione, dwork, ldwork,
work, W(i+1,i), stream );
#ifdef PROFILE_SYMV
magma_setdevice(0);
magma_event_record(stop, stream[0][0]);
#endif
trace_cpu_start( 0, "gemv", "gemv" );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
trace_cpu_end( 0 );
/* overlap update */
if( i > 0 && i+1 < n )
{
magma_int_t ip1 = i+1;
trace_cpu_start( 0, "gemv", "gemv" );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(ip1, 0), &ldw);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(ip1, 0), &lda,
W(ip1, 0), &ldw, &c_one, A(ip1, ip1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(ip1, 0), &ldw);
lapackf77_zlacgv(&i, A(ip1 ,0), &lda);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(ip1, 0), &ldw,
A(ip1, 0), &lda, &c_one, A(ip1, ip1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, A(ip1, 0), &lda);
#endif
trace_cpu_end( 0 );
}
/* synchronize */
magmablas_zhemv_sync(num_gpus, k, i_n, work, W(i+1,i), stream );
#ifdef PROFILE_SYMV
cudaEventElapsedTime(&etime, start, stop);
mv_time += (etime/1000.0);
times[1+(i_n/(n0/10))] += (etime/1000.0);
#endif
trace_cpu_start( 0, "axpy", "axpy" );
if (i!=0)
blasf77_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("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_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i]* -.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
trace_cpu_end( 0 );
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
if( k > 1 ) {
magma_zsetvector_async( n, W(0,i), 1, dW(id, 0, i), 1, stream[id][1] );
} else {
magma_zsetvector_async( n, W(0,i), 1, dW(id, 0, i), 1, stream[id][0] );
}
}
}
}
}
#ifdef PROFILE_SYMV
magma_setdevice(0);
magma_event_destory( start );
magma_event_destory( stop );
#endif
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
if( k > 1) magma_queue_sync(stream[id][1]);
}
free(f);
return mv_time;
} /* zlatrd_ */