// 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,
double *A, magma_int_t lda, magma_int_t *ipiv,
double *x, magma_int_t ldx,
double *b, magma_int_t ldb)
{
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_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_dgemv( "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_dlange( MagmaFullStr, &n, &n, A, &lda, work );
norm_r = lapackf77_dlange( MagmaFullStr, &n, &ione, b, &n, work );
norm_x = lapackf77_dlange( MagmaFullStr, &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_dprint( 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,
double *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 double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_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;
double *x, *b;
// initialize RHS
TESTING_MALLOC_CPU( x, double, n );
TESTING_MALLOC_CPU( b, double, n );
lapackf77_dlarnv( &ione, ISEED, &n, b );
blasf77_dcopy( &n, b, &ione, x, &ione );
// solve Ax = b
lapackf77_dgetrs( "Notrans", &n, &ione, A, &lda, ipiv, x, &n, &info );
if (info != 0)
printf("lapackf77_dgetrs 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_dgemv( "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_dlange( "F", &m, &n, A, &lda, work );
norm_r = lapackf77_dlange( "F", &n, &ione, b, &n, work );
norm_x = lapackf77_dlange( "F", &n, &ione, x, &n, work );
//printf( "r=\n" ); magma_dprint( 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);
}
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;
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 1.5, -2.3 );
double beta = MAGMA_D_MAKE( -0.6, 0.8 );
double *A, *X, *Y, *Ycublas, *Ymagma;
double *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_DGEMV( 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, double, sizeA );
TESTING_MALLOC_CPU( X, double, sizeX );
TESTING_MALLOC_CPU( Y, double, sizeY );
TESTING_MALLOC_CPU( Ycublas, double, sizeY );
TESTING_MALLOC_CPU( Ymagma, double, sizeY );
TESTING_MALLOC_DEV( dA, double, sizeA );
TESTING_MALLOC_DEV( dX, double, sizeX );
TESTING_MALLOC_DEV( dY, double, sizeY );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &sizeA, A );
lapackf77_dlarnv( &ione, ISEED, &sizeX, X );
lapackf77_dlarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_dsetmatrix( M, N, A, lda, dA, lda );
magma_dsetvector( Xm, X, incx, dX, incx );
magma_dsetvector( Ym, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasDgemv( 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_dgetvector( Ym, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_dsetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_dgemv( 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_dgetvector( Ym, dY, incx, Ymagma, incx );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_dgemv( 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_daxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_dlange( "M", &Ym, &ione, Ymagma, &Ym, work ) / Ym;
blasf77_daxpy( &Ym, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_dlange( "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;
}
/**
Purpose
-------
DLAHR2 reduces the first NB columns of a real 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 DGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
A DOUBLE_PRECISION 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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T DOUBLE_PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y DOUBLE_PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
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 real scalar, and v is a real 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:
@verbatim
( 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 )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_dgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
double *A, magma_int_t lda,
double *tau,
double *T, magma_int_t ldt,
double *Y, magma_int_t ldy,
struct dgehrd_data* data )
{
#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)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
double 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;
double scale;
magma_int_t i;
double ei = MAGMA_D_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_dlaset( MagmaFull, nb, nb, c_zero, c_zero, dV(d,k,0), ldv );
}
// set all Y=0
lapackf77_dlaset( "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_D_NEGATE( tau[i-1] );
blasf77_daxpy( &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_dcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_dtrmv( "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_dgemv( "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_dtrmv( "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_dgemv( "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_dtrmv( "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_daxpy( &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_dlarfg( &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_dsetvector_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_dlacpy( MagmaFull, 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_dgemv( MagmaNoTrans, 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_dgetvector_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_D_NEGATE( tau[i] );
blasf77_dgemv( "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_dtrmv( "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 real, conjugate row of V, and undo afterwards
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &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_dgemv( "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_dlacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_dgemv( "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_daxpy( &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_dgemm( MagmaNoTrans, MagmaNoTrans, 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_dgetmatrix_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_daxpy( &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_dsetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->streams[d] );
magma_dsetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->streams[d] );
}
return 0;
} /* magma_dlahr2 */
extern "C" magma_int_t
magma_dlabrd_gpu( magma_int_t m, magma_int_t n, magma_int_t nb,
double *a, magma_int_t lda, double *da, magma_int_t ldda,
double *d, double *e, double *tauq, double *taup,
double *x, magma_int_t ldx, double *dx, magma_int_t lddx,
double *y, magma_int_t ldy, double *dy, magma_int_t lddy)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DLABRD reduces the first NB rows and columns of a real 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) DOUBLE_PRECISION 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) DOUBLE_PRECISION array, dimension (NB)
The diagonal elements of the first NB rows and columns of
the reduced matrix. D(i) = A(i,i).
E (output) DOUBLE_PRECISION array, dimension (NB)
The off-diagonal elements of the first NB rows and columns of
the reduced matrix.
TAUQ (output) DOUBLE_PRECISION array dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
TAUP (output) DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
X (output) DOUBLE_PRECISION 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) DOUBLE_PRECISION 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 real scalars, and v and u are real 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).
===================================================================== */
/* Table of constant values */
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
magma_int_t c__1 = 1;
/* System generated locals */
magma_int_t a_dim1, a_offset, x_dim1, x_offset, y_dim1, y_offset, i__2, i__3;
/* Local variables */
magma_int_t i__;
double 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-= 1 + lddx;
y_dim1 = ldy;
y_offset = 1 + y_dim1;
y -= y_offset;
dy-= 1 + lddy;
/* Function Body */
if (m <= 0 || n <= 0) {
return 0;
}
double *f;
magma_queue_t stream;
magma_queue_create( &stream );
magma_dmalloc_cpu( &f, max(n,m) );
assert( f != NULL ); // TODO return error, or allocate outside dlatrd
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_dlacgv( &i__3, &y[i__+y_dim1], &ldy );
#endif
blasf77_dgemv("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_dlacgv( &i__3, &y[i__+y_dim1], &ldy );
#endif
blasf77_dgemv("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_dlarfg(&i__2, &alpha,
&a[min(i__3,m) + i__ * a_dim1], &c__1, &tauq[i__]);
d[i__] = MAGMA_D_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_dsetvector( i__2,
a + i__ + i__ * a_dim1, 1,
da+(i__-1)+(i__-1)* (ldda), 1 );
// 2. Multiply ---------------------------------------------
magma_dgemv(MagmaTrans, i__2, i__3, c_one,
da + (i__-1) + ((i__-1) + 1) * (ldda), ldda,
da + (i__-1) + (i__-1) * (ldda), c__1, c_zero,
dy + i__ + 1 + i__ * y_dim1, c__1);
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__3, 1,
dy+i__+1+i__*y_dim1, y_dim1,
y+i__+1+i__*y_dim1, y_dim1, stream );
i__2 = m - i__ + 1;
i__3 = i__ - 1;
blasf77_dgemv(MagmaTransStr, &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_dgemv("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_dgemv(MagmaTransStr, &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_queue_sync( stream );
if (i__3!=0){
i__2 = n - i__;
blasf77_daxpy(&i__2, &c_one, f,&c__1, &y[i__+1+i__*y_dim1],&c__1);
}
i__2 = i__ - 1;
i__3 = n - i__;
blasf77_dgemv(MagmaTransStr, &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_dscal(&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_dlacgv( &i__2, &a[i__+(i__+1)*a_dim1], &lda );
lapackf77_dlacgv( &i__, &a[i__+a_dim1], &lda );
#endif
blasf77_dgemv("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_dlacgv( &i__, &a[i__+a_dim1], &lda );
lapackf77_dlacgv( &i__2, &x[i__+x_dim1], &ldx );
#endif
blasf77_dgemv(MagmaTransStr, &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_dlacgv( &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_dlarfg(&i__2, &alpha, &a[i__ + min(
i__3,n) * a_dim1], &lda, &taup[i__]);
e[i__] = MAGMA_D_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_dsetvector( i__3,
a + i__ + (i__ +1)* a_dim1, lda,
da+(i__-1)+((i__-1)+1)*(ldda), ldda );
// 2. Multiply ---------------------------------------------
//magma_dcopy(i__3, da+(i__-1)+((i__-1)+1)*(ldda), ldda,
// dy + 1 + lddy, 1);
magma_dgemv(MagmaNoTrans, i__2, i__3, c_one,
da + (i__-1)+1+ ((i__-1)+1) * (ldda), ldda,
da + (i__-1) + ((i__-1)+1) * (ldda), ldda,
//dy + 1 + lddy, 1,
c_zero, dx + i__ + 1 + i__ * x_dim1, c__1);
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__2, 1,
dx+i__+1+i__*x_dim1, x_dim1,
x+i__+1+i__*x_dim1, x_dim1, stream );
i__2 = n - i__;
blasf77_dgemv(MagmaTransStr, &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_dgemv("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_dgemv("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_queue_sync( stream );
if (i__!=0){
i__2 = m - i__;
blasf77_daxpy(&i__2, &c_one, f,&c__1, &x[i__+1+i__*x_dim1],&c__1);
}
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_dgemv("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_dscal(&i__2, &taup[i__], &x[i__ + 1 + i__ * x_dim1], &c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
i__2 = n - i__;
lapackf77_dlacgv( &i__2, &a[i__+(i__+1)*a_dim1], &lda );
// 4. Send the block reflector A(i+1:m,i) to the GPU after DLACGV()
magma_dsetvector( i__2,
a + i__ + (i__ +1)* a_dim1, lda,
da+(i__-1)+((i__-1)+1)*(ldda), ldda );
#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_dlacgv(&i__2, &a[i__ + i__ * a_dim1], &lda);
lapackf77_dlacgv(&i__3, &a[i__ + a_dim1], &lda);
#endif
blasf77_dgemv("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_dlacgv(&i__3, &a[i__ + a_dim1], &lda);
lapackf77_dlacgv(&i__3, &x[i__ + x_dim1], &ldx);
#endif
i__3 = n - i__ + 1;
blasf77_dgemv(MagmaTransStr, &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_dlacgv(&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_dlarfg(&i__2, &alpha,
&a[i__ + min(i__3,n) * a_dim1], &lda, &taup[i__]);
d[i__] = MAGMA_D_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_dsetvector( i__3,
a + i__ + i__ * a_dim1, lda,
da+(i__-1)+(i__-1)* (ldda), ldda );
// 2. Multiply ---------------------------------------------
//magma_dcopy(i__3, da+(i__-1)+(i__-1)*(ldda), ldda,
// dy + 1 + lddy, 1);
magma_dgemv(MagmaNoTrans, i__2, i__3, c_one,
da + (i__-1)+1 + (i__-1) * ldda, ldda,
da + (i__-1) + (i__-1) * ldda, ldda,
// dy + 1 + lddy, 1,
c_zero,
dx + i__ + 1 + i__ * x_dim1, c__1);
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__2, 1,
dx+i__+1+i__*x_dim1, x_dim1,
x+i__+1+i__*x_dim1, x_dim1, stream );
i__2 = n - i__ + 1;
i__3 = i__ - 1;
blasf77_dgemv(MagmaTransStr, &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_dgemv("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_dgemv("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_queue_sync( stream );
if (i__2!=0){
i__3 = m - i__;
blasf77_daxpy(&i__3, &c_one, f,&c__1, &x[i__+1+i__*x_dim1],&c__1);
}
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_dgemv("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_dscal(&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_dlacgv(&i__2, &a[i__ + i__ * a_dim1], &lda);
magma_dsetvector( i__2,
a + i__ + (i__ )* a_dim1, lda,
da+(i__-1)+ (i__-1)*(ldda), ldda );
#endif
/* Update A(i+1:m,i) */
i__2 = m - i__;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv(&i__3, &y[i__ + y_dim1], &ldy);
#endif
blasf77_dgemv("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_dlacgv(&i__3, &y[i__ + y_dim1], &ldy);
#endif
blasf77_dgemv("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_dlarfg(&i__2, &alpha,
&a[min(i__3,m) + i__ * a_dim1], &c__1, &tauq[i__]);
e[i__] = MAGMA_D_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_dsetvector( i__2,
a + i__ +1+ i__ * a_dim1, 1,
da+(i__-1)+1+ (i__-1)*(ldda), 1 );
// 2. Multiply ---------------------------------------------
magma_dgemv(MagmaTrans, i__2, i__3, c_one,
da + (i__-1)+1+ ((i__-1)+1) * ldda, ldda,
da + (i__-1)+1+ (i__-1) * ldda, c__1,
c_zero, dy + i__ + 1 + i__ * y_dim1, c__1);
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__3, 1,
dy+i__+1+i__*y_dim1, y_dim1,
y+i__+1+i__*y_dim1, y_dim1, stream );
i__2 = m - i__;
i__3 = i__ - 1;
blasf77_dgemv(MagmaTransStr, &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_dgemv("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_dgemv(MagmaTransStr, &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_queue_sync( stream );
if (i__3!=0){
i__2 = n - i__;
blasf77_daxpy(&i__2, &c_one, f,&c__1, &y[i__+1+i__*y_dim1],&c__1);
}
i__2 = n - i__;
blasf77_dgemv(MagmaTransStr, &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_dscal(&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_dlacgv(&i__2, &a[i__ + i__ * a_dim1], &lda);
magma_dsetvector( i__2,
a + i__ + (i__ )* a_dim1, lda,
da+(i__-1)+ (i__-1)*(ldda), ldda );
}
#endif
}
}
magma_queue_destroy( stream );
magma_free_cpu(f);
return MAGMA_SUCCESS;
} /* dlabrd */
/**
Purpose
-------
DLABRD reduces the first NB rows and columns of a real 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 DGEBRD.
Arguments
---------
@param[in]
m INTEGER
The number of rows in the matrix A.
@param[in]
n INTEGER
The number of columns in the matrix A.
@param[in]
nb INTEGER
The number of leading rows and columns of A to be reduced.
@param[in,out]
A DOUBLE_PRECISION 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.
\n
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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
dA DOUBLE_PRECISION array, dimension (LDDA,N)
Copy of A on GPU.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
d DOUBLE_PRECISION array, dimension (NB)
The diagonal elements of the first NB rows and columns of
the reduced matrix. D(i) = A(i,i).
@param[out]
e DOUBLE_PRECISION array, dimension (NB)
The off-diagonal elements of the first NB rows and columns of
the reduced matrix.
@param[out]
tauq DOUBLE_PRECISION array dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
@param[out]
taup DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
@param[out]
X DOUBLE_PRECISION array, dimension (LDX,NB)
The m-by-nb matrix X required to update the unreduced part
of A.
@param[in]
ldx INTEGER
The leading dimension of the array X. LDX >= M.
@param[out]
dX DOUBLE_PRECISION array, dimension (LDDX,NB)
Copy of X on GPU.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= M.
@param[out]
Y DOUBLE_PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y required to update the unreduced part
of A.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[out]
dY DOUBLE_PRECISION array, dimension (LDDY,NB)
Copy of Y on GPU.
@param[in]
lddy INTEGER
The leading dimension of the array dY. LDDY >= 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 real scalars, and v and u are real 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:
@verbatim
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 )
@endverbatim
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).
@ingroup magma_dgesvd_aux
********************************************************************/
extern "C" magma_int_t
magma_dlabrd_gpu( magma_int_t m, magma_int_t n, magma_int_t nb,
double *A, magma_int_t lda,
double *dA, magma_int_t ldda,
double *d, double *e, double *tauq, double *taup,
double *X, magma_int_t ldx,
double *dX, magma_int_t lddx,
double *Y, magma_int_t ldy,
double *dY, magma_int_t lddy)
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
#define X(i_,j_) (X + (i_) + (j_)*ldx)
#define Y(i_,j_) (Y + (i_) + (j_)*ldy)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dY(i_,j_) (dY + (i_) + (j_)*lddy)
#define dX(i_,j_) (dX + (i_) + (j_)*lddx)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
magma_int_t ione = 1;
magma_int_t i__2, i__3;
magma_int_t i;
double alpha;
A -= 1 + lda;
X -= 1 + ldx;
dX -= 1 + lddx;
Y -= 1 + ldy;
dY -= 1 + lddy;
--d;
--e;
--tauq;
--taup;
/* Quick return if possible */
magma_int_t info = 0;
if (m <= 0 || n <= 0) {
return info;
}
double *f;
magma_queue_t stream;
magma_queue_create( &stream );
magma_dmalloc_cpu( &f, max(n,m) );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
return info;
}
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_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i,1), &lda,
Y(i,1), &ldy, &c_one,
A(i,i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i,1), &ldx,
A(1,i), &ione, &c_one,
A(i,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+1:m,i) */
alpha = *A(i,i);
i__2 = m - i + 1;
i__3 = i + 1;
lapackf77_dlarfg( &i__2, &alpha, A(min(i__3,m),i), &ione, &tauq[i] );
d[i] = MAGMA_D_REAL( alpha );
if (i < n) {
*A(i,i) = 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_dsetvector( i__2,
A(i,i), 1,
dA(i-1,i-1), 1 );
// 2. Multiply ---------------------------------------------
magma_dgemv( MagmaConjTrans, i__2, i__3, c_one,
dA(i-1,i), ldda,
dA(i-1,i-1), ione, c_zero,
dY(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__3, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, stream );
i__2 = m - i + 1;
i__3 = i - 1;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
A(i,1), &lda,
A(i,i), &ione, &c_zero,
Y(1,i), &ione );
i__2 = n - i;
i__3 = i - 1;
blasf77_dgemv( "N", &i__2, &i__3, &c_neg_one,
Y(i+1,1), &ldy,
Y(1,i), &ione, &c_zero,
f, &ione );
i__2 = m - i + 1;
i__3 = i - 1;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
X(i,1), &ldx,
A(i,i), &ione, &c_zero,
Y(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__3 != 0) {
i__2 = n - i;
blasf77_daxpy( &i__2, &c_one, f, &ione, Y(i+1,i), &ione );
}
i__2 = i - 1;
i__3 = n - i;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
A(1,i+1), &lda,
Y(1,i), &ione, &c_one,
Y(i+1,i), &ione );
i__2 = n - i;
blasf77_dscal( &i__2, &tauq[i], Y(i+1,i), &ione );
/* Update A(i,i+1:n) */
i__2 = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, A(i,i+1), &lda );
lapackf77_dlacgv( &i, A(i,1), &lda );
#endif
blasf77_dgemv( "No transpose", &i__2, &i, &c_neg_one,
Y(i+1,1), &ldy,
A(i,1), &lda, &c_one,
A(i,i+1), &lda );
i__2 = i - 1;
i__3 = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i, A(i,1), &lda );
lapackf77_dlacgv( &i__2, X(i,1), &ldx );
#endif
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
A(1,i+1), &lda,
X(i,1), &ldx, &c_one,
A(i,i+1), &lda );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, X(i,1), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+2:n) */
i__2 = n - i;
i__3 = i + 2;
alpha = *A(i,i+1);
lapackf77_dlarfg( &i__2, &alpha, A(i,min(i__3,n)), &lda, &taup[i] );
e[i] = MAGMA_D_REAL( alpha );
*A(i,i+1) = 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_dsetvector( i__3,
A(i,i+1), lda,
dA(i-1,i), ldda );
// 2. Multiply ---------------------------------------------
//magma_dcopy( i__3, dA(i-1,i), ldda, dY(1,1), 1 );
magma_dgemv( MagmaNoTrans, i__2, i__3, c_one,
dA(i,i), ldda,
dA(i-1,i), ldda,
//dY(1,1), 1,
c_zero,
dX(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__2, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, stream );
i__2 = n - i;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i, &c_one,
Y(i+1,1), &ldy,
A(i,i+1), &lda, &c_zero,
X(1,i), &ione );
i__2 = m - i;
blasf77_dgemv( "N", &i__2, &i, &c_neg_one,
A(i+1,1), &lda,
X(1,i), &ione, &c_zero,
f, &ione );
i__2 = i - 1;
i__3 = n - i;
blasf77_dgemv( "N", &i__2, &i__3, &c_one,
A(1,i+1), &lda,
A(i,i+1), &lda, &c_zero,
X(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i != 0) {
i__2 = m - i;
blasf77_daxpy( &i__2, &c_one, f, &ione, X(i+1,i), &ione );
}
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i+1,1), &ldx,
X(1,i), &ione, &c_one,
X(i+1,i), &ione );
i__2 = m - i;
blasf77_dscal( &i__2, &taup[i], X(i+1,i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
i__2 = n - i;
lapackf77_dlacgv( &i__2, A(i,i+1), &lda );
// 4. Send the block reflector A(i+1:m,i) to the GPU after DLACGV()
magma_dsetvector( i__2,
A(i,i+1), lda,
dA(i-1,i), ldda );
#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_dlacgv( &i__2, A(i,i), &lda );
lapackf77_dlacgv( &i__3, A(i,1), &lda );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
Y(i,1), &ldy,
A(i,1), &lda, &c_one,
A(i,i), &lda );
i__2 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, A(i,1), &lda );
lapackf77_dlacgv( &i__3, X(i,1), &ldx );
#endif
i__3 = n - i + 1;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
A(1,i), &lda,
X(i,1), &ldx, &c_one,
A(i,i), &lda );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, X(i,1), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+1:n) */
i__2 = n - i + 1;
i__3 = i + 1;
alpha = *A(i,i);
lapackf77_dlarfg( &i__2, &alpha, A(i,min(i__3,n)), &lda, &taup[i] );
d[i] = MAGMA_D_REAL( alpha );
if (i < m) {
*A(i,i) = 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_dsetvector( i__3,
A(i,i), lda,
dA(i-1,i-1), ldda );
// 2. Multiply ---------------------------------------------
//magma_dcopy( i__3, dA(i-1,i-1), ldda, dY(1,1), 1 );
magma_dgemv( MagmaNoTrans, i__2, i__3, c_one,
dA(i,i-1), ldda,
dA(i-1,i-1), ldda,
//dY(1,1), 1,
c_zero,
dX(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__2, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, stream );
i__2 = n - i + 1;
i__3 = i - 1;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
Y(i,1), &ldy,
A(i,i), &lda, &c_zero,
X(1,i), &ione );
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i+1,1), &lda,
X(1,i), &ione, &c_zero,
f, &ione );
i__2 = i - 1;
i__3 = n - i + 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_one,
A(1,i), &lda,
A(i,i), &lda, &c_zero,
X(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__2 != 0) {
i__3 = m - i;
blasf77_daxpy( &i__3, &c_one, f, &ione, X(i+1,i), &ione );
}
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i+1,1), &ldx,
X(1,i), &ione, &c_one,
X(i+1,i), &ione );
i__2 = m - i;
blasf77_dscal( &i__2, &taup[i], X(i+1,i), &ione );
i__2 = n - i + 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__2, A(i,i), &lda );
magma_dsetvector( i__2,
A(i,i), lda,
dA(i-1,i-1), ldda );
#endif
/* Update A(i+1:m,i) */
i__2 = m - i;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i+1,1), &lda,
Y(i,1), &ldy, &c_one,
A(i+1,i), &ione );
i__2 = m - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_dgemv( "No transpose", &i__2, &i, &c_neg_one,
X(i+1,1), &ldx,
A(1,i), &ione, &c_one,
A(i+1,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+2:m,i) */
i__2 = m - i;
i__3 = i + 2;
alpha = *A(i+1,i);
lapackf77_dlarfg( &i__2, &alpha, A(min(i__3,m),i), &ione, &tauq[i] );
e[i] = MAGMA_D_REAL( alpha );
*A(i+1,i) = 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_dsetvector( i__2,
A(i+1,i), 1,
dA(i,i-1), 1 );
// 2. Multiply ---------------------------------------------
magma_dgemv( MagmaConjTrans, i__2, i__3, c_one,
dA(i,i), ldda,
dA(i,i-1), ione, c_zero,
dY(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_dgetmatrix_async( i__3, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, stream );
i__2 = m - i;
i__3 = i - 1;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
A(i+1,1), &lda,
A(i+1,i), &ione, &c_zero,
Y(1,i), &ione );
i__2 = n - i;
i__3 = i - 1;
blasf77_dgemv( "No transpose", &i__2, &i__3, &c_neg_one,
Y(i+1,1), &ldy,
Y(1,i), &ione, &c_zero,
f, &ione );
i__2 = m - i;
blasf77_dgemv( MagmaConjTransStr, &i__2, &i, &c_one,
X(i+1,1), &ldx,
A(i+1,i), &ione, &c_zero,
Y(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__3 != 0) {
i__2 = n - i;
blasf77_daxpy( &i__2, &c_one, f, &ione, Y(i+1,i), &ione );
}
i__2 = n - i;
blasf77_dgemv( MagmaConjTransStr, &i, &i__2, &c_neg_one,
A(1,i+1), &lda,
Y(1,i), &ione, &c_one,
Y(i+1,i), &ione );
i__2 = n - i;
blasf77_dscal( &i__2, &tauq[i], Y(i+1,i), &ione );
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else {
i__2 = n - i + 1;
lapackf77_dlacgv( &i__2, A(i,i), &lda );
magma_dsetvector( i__2,
A(i,i), lda,
dA(i-1,i-1), ldda );
}
#endif
}
}
magma_queue_destroy( stream );
magma_free_cpu( f );
return info;
} /* magma_dlabrd_gpu */
extern "C" magma_int_t
magma_dlatrd2(char uplo, magma_int_t n, magma_int_t nb,
double *a, magma_int_t lda,
double *e, double *tau,
double *w, magma_int_t ldw,
double *da, magma_int_t ldda,
double *dw, magma_int_t lddw,
double *dwork, magma_int_t ldwork)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
DLATRD2 reduces NB rows and columns of a real symmetric matrix A to
symmetric 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', DLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = 'L', DLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by DSYTRD2_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
symmetric 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) DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric 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) DOUBLE_PRECISION 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) DOUBLE_PRECISION 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) DOUBLE_PRECISION 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 real scalar, and v is a real 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 real scalar, and v is a real 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 symmetric 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;
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
double value = MAGMA_D_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
double alpha;
double *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_dmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside dlatrd
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_dlacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_dgemv("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_dlacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_dlacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_dgemv("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_dlacgv(&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_dlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_D_REAL( alpha );
*A(i-1,i) = MAGMA_D_MAKE( 1, 0 );
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_dsetvector( i, A(0, i), 1, dA(0, i), 1 );
//#if (GPUSHMEM < 200)
//magma_dsymv(MagmaUpper, i, c_one, dA(0, 0), ldda,
// dA(0, i), ione, c_zero, dW(0, iw), ione);
//#else
magmablas_dsymv_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_dgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_dgemv(MagmaTransStr, &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_dgemv("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_dgemv(MagmaTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_dgemv("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_dscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_ddot_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_ddot( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_daxpy(&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_dlacgv(&i, W(i, 0), &ldw);
#endif
blasf77_dgemv("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_dlacgv(&i, W(i, 0), &ldw);
lapackf77_dlacgv(&i, A(i ,0), &lda);
#endif
blasf77_dgemv("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_dlacgv(&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_dlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_D_REAL( alpha );
*A(i+1,i) = MAGMA_D_MAKE( 1, 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_dsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
//#if (GPUSHMEM < 200)
//magma_dsymv(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_dsymv_work('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_dgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_dgemv(MagmaTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_dgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_dgemv(MagmaTransStr, &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_daxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_dgemv("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_dscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_ddot_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_ddot( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_daxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* dlatrd */
/**
Purpose
-------
DLAHR2 reduces the first NB columns of a real 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 DGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
dA DOUBLE PRECISION array on the GPU, dimension (LDDA,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.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[out]
dV DOUBLE PRECISION array on the GPU, dimension (LDDV, NB)
On exit this n-by-nb array contains the Householder vectors of the transformation.
@param[in]
lddv INTEGER
The leading dimension of the array dV. LDDV >= max(1,N).
@param[in,out]
A DOUBLE PRECISION 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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T DOUBLE PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y DOUBLE PRECISION array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in]
queue magma_queue_t
Queue to execute in.
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 real scalar, and v is a real 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:
@verbatim
( 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 )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_dgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlahr2(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dV, magma_int_t lddv,
double *A, magma_int_t lda,
double *tau,
double *T, magma_int_t ldt,
double *Y, magma_int_t ldy,
magma_queue_t queue )
{
#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(i_,j_) (dA + (i_) + (j_)*ldda)
#define dV(i_,j_) (dV + (i_) + (j_)*lddv)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
magma_int_t n_k_i_1, n_k;
double scale;
magma_int_t i;
double ei = MAGMA_D_ZERO;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (k < 0 || k > n) {
info = -2;
} else if (nb < 1 || nb > n) {
info = -3;
} else if (ldda < max(1,n)) {
info = -5;
} else if (lddv < max(1,n)) {
info = -7;
} else if (lda < max(1,n)) {
info = -9;
} else if (ldt < max(1,nb)) {
info = -12;
} else if (ldy < max(1,n)) {
info = -13;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
// adjust from 1-based indexing
k -= 1;
if (n <= 1)
return info;
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Update A(k:n-1,i); Update i-th column of A - Y * T * V'
// This updates one more row than LAPACK does (row k),
// making the block above the panel an even multiple of nb.
// Use last column of T as workspace, w.
// w(0:i-1, nb-1) = VA(k+i, 0:i-1)'
blasf77_dcopy( &i,
A(k+i,0), &lda,
T(0,nb-1), &ione );
#ifdef COMPLEX
// If real, conjugate row of V.
lapackf77_dlacgv(&i, T(0,nb-1), &ione);
#endif
// w = T(0:i-1, 0:i-1) * w
blasf77_dtrmv( "Upper", "No trans", "No trans", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// A(k:n-1, i) -= Y(k:n-1, 0:i-1) * w
blasf77_dgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, 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_dcopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_dtrmv( "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_dgemv( "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_dtrmv( "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_dgemv( "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_dtrmv( "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_daxpy( &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_dlarfg( &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;
// dV(i+1:n-k-1, i) = VA(k+i+1:n-1, i)
magma_dsetvector( n_k_i_1,
A(k+i+1,i), 1,
dV(i+1,i), 1, queue );
// Compute Y(k+1:n,i) = A vi
// dA(k:n-1, i) = dA(k:n-1, i+1:n-k-1) * dV(i+1:n-k-1, i)
magma_dgemv( MagmaNoTrans, n_k, n_k_i_1,
c_one, dA(k,i+1), ldda,
dV(i+1,i), ione,
c_zero, dA(k,i), ione, queue );
// 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_D_NEGATE( tau[i]);
blasf77_dgemv( "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_dtrmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// Y(k:n-1, i) = dA(k:n-1, i)
magma_dgetvector( n-k,
dA(k,i), 1,
Y(k,i), 1, queue );
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
return info;
} /* magma_dlahr2 */
/**
Purpose
-------
DLATRD2 reduces NB rows and columns of a real symmetric matrix A to
symmetric 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, DLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, DLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by DSYTRD2_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
symmetric 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 DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric 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 DOUBLE_PRECISION 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 DOUBLE_PRECISION 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 DOUBLE_PRECISION 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, 2*nb) ??
@param
lddw TODO: lddw >= n ??
@param
dwork TODO: dimension (ldwork) ??
@param
ldwork TODO: ldwork >= ceil(n/64)*ldda ??
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 real scalar, and v is a real 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 real scalar, and v is a real 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 symmetric 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_dsyev_aux
********************************************************************/
extern "C" magma_int_t
magma_dlatrd2(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
double *A, magma_int_t lda,
double *e, double *tau,
double *W, magma_int_t ldw,
magmaDouble_ptr dA, magma_int_t ldda,
magmaDouble_ptr dW, magma_int_t lddw,
magmaDouble_ptr dwork, magma_int_t ldwork)
{
#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)
const double c_neg_one = MAGMA_D_NEG_ONE;
const double c_one = MAGMA_D_ONE;
const double c_zero = MAGMA_D_ZERO;
const magma_int_t ione = 1;
double 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;
} else if ( ldwork < ldda*ceildiv(n,64) ) {
info = -15;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Quick return if possible */
if (n == 0) {
return info;
}
magma_queue_t stream;
magma_queue_create( &stream );
double *f;
magma_dmalloc_cpu( &f, n );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
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) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv( &i_n, W(i, iw+1), &ldw );
#endif
blasf77_dgemv( "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_dlacgv( &i_n, W(i, iw+1), &ldw );
lapackf77_dlacgv( &i_n, A(i, i+1), &lda );
#endif
blasf77_dgemv( "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_dlacgv( &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_dlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_D_REAL( alpha );
*A(i-1,i) = MAGMA_D_ONE;
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_dsetvector_async( i, A(0, i), 1, dA(0, i), 1, stream );
magmablas_dsymv_work( MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork, stream );
// 2. Start getting the result back (asynchronously)
magma_dgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw), ldw, stream );
if (i < n-1) {
blasf77_dgemv( MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
}
// 3. Here we need dsymv result W(0, iw)
magma_queue_sync( stream );
if (i < n-1) {
blasf77_dgemv( "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_dgemv( MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
blasf77_dgemv( "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_dscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_ddot( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_daxpy( &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_dlacgv( &i, W(i, 0), &ldw );
#endif
blasf77_dgemv( "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_dlacgv( &i, W(i, 0), &ldw );
lapackf77_dlacgv( &i, A(i, 0), &lda );
#endif
blasf77_dgemv( "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_dlacgv( &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_dlarfg( &i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_D_REAL( alpha );
*A(i+1,i) = MAGMA_D_ONE;
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_dsetvector_async( i_n, A(i+1, i), 1, dA(i+1, i), 1, stream );
magmablas_dsymv_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, stream );
// 2. Start getting the result back (asynchronously)
magma_dgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_dgemv( MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
blasf77_dgemv( "No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione );
blasf77_dgemv( MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
// 3. Here we need dsymv result W(i+1, i)
magma_queue_sync( stream );
if (i != 0)
blasf77_daxpy( &i_n, &c_one, f, &ione, W(i+1, i), &ione );
blasf77_dgemv( "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_dscal( &i_n, &tau[i], W(i+1,i), &ione );
value = magma_cblas_ddot( i_n, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_daxpy( &i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
}
}
}
magma_free_cpu( f );
magma_queue_destroy( stream );
return info;
} /* magma_dlatrd */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *d_A, *h_R, *tau, *dT, *h_work, tmp[1];
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol;
opts.lapack |= (opts.version == 2 && opts.check == 2); // check (-c2) implies lapack (-l)
if ( opts.version != 2 && opts.check == 1 ) {
printf( " ===================================================================\n"
" NOTE: -c check for this version will be wrong\n"
" because tester ignores the special structure of MAGMA dgeqrf resuls.\n"
" We reset it to -c2.\n"
" ===================================================================\n\n");
opts.check = 2;
}
if ( opts.version == 2 ) {
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1*eps) ||I-Q'Q||_1 / (M*eps)\n");
printf("=========================================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F / ||A||_F\n");
printf("=======================================================================\n");
}
tol = 1.0;
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||Ax-b||_F/(N*||A||_F*||x||_F)\n");
printf("====================================================================================\n");
tol = opts.tolerance * lapackf77_dlamch("E");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_work, double, lwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
if ( opts.version == 2 ) {
magma_dgeqrf2_gpu( M, N, d_A, ldda, tau, &info);
}
else {
nb = magma_get_dgeqrf_nb( M );
size = (2*min(M, N) + (N+31)/32*32 )*nb;
TESTING_MALLOC_DEV( dT, double, size );
if ( opts.version == 3 ) {
magma_dgeqrf3_gpu( M, N, d_A, ldda, tau, dT, &info);
}
else {
magma_dgeqrf_gpu( M, N, d_A, ldda, tau, dT, &info);
}
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
double *tau2;
TESTING_MALLOC_CPU( tau2, double, min_mn );
cpu_time = magma_wtime();
lapackf77_dgeqrf(&M, &N, h_A, &lda, tau2, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( tau2 );
}
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
double *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
magma_dgetmatrix( M, N, d_A, ldda, h_R, M );
TESTING_MALLOC_CPU( h_W1, double, n2 ); // Q
TESTING_MALLOC_CPU( h_W2, double, n2 ); // R
TESTING_MALLOC_CPU( h_W3, double, lwork ); // WORK
TESTING_MALLOC_CPU( h_RW, double, M ); // RWORK
lapackf77_dlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_dqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0], results[1] );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0], results[1] );
}
// todo also check results[1] < tol?
printf(" %s\n", (results[0] < tol ? "ok" : "failed"));
status += ! (results[0] < tol);
TESTING_FREE_CPU( h_W1 );
TESTING_FREE_CPU( h_W2 );
TESTING_FREE_CPU( h_W3 );
TESTING_FREE_CPU( h_RW );
}
else if ( opts.check == 2 ) {
if ( opts.version == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_dgetmatrix( M, N, d_A, ldda, h_R, M );
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work) / error;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf(" %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else if ( M >= N ) {
magma_int_t lwork;
double *x, *b, *d_B, *hwork;
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const magma_int_t ione = 1;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, double, N );
TESTING_MALLOC_CPU( b, double, M );
lapackf77_dlarnv( &ione, ISEED, &N, x );
blasf77_dgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, double, M );
magma_dsetvector( M, b, 1, d_B, 1 );
if ( opts.version == 1 ) {
// allocate hwork
magma_dgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, double, lwork );
// solve linear system
magma_dgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork, &info );
if (info != 0)
printf("magma_dgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
else {
// allocate hwork
magma_dgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, double, lwork );
// solve linear system
magma_dgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork, &info );
if (info != 0)
printf("magma_dgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
magma_dgetvector( N, d_B, 1, x, 1 );
// compute r = Ax - b, saved in b
lapackf77_dlarnv( &ione, ISEED2, &n2, h_A );
blasf77_dgemv( "Notrans", &M, &N, &c_one, h_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_dlange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_dlange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_dlange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (N * norm_A * norm_x);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, error );
}
printf(" %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- ",
(int) M, (int) N, gpu_perf, gpu_time);
}
printf("%s\n", (opts.check != 0 ? " (error check only for M >= N)" : ""));
}
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version != 2 )
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const double d_neg_one = MAGMA_D_NEG_ONE;
const double d_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const double c_one = MAGMA_D_ONE;
const double c_zero = MAGMA_D_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double Anorm, error=0, error2=0;
double *h_A, *h_R, *tau, *h_work, tmp[1];
magmaDouble_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf( "version %d\n", (int) opts.version );
if ( opts.version == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |R - Q^H*A| |I - Q^H*Q|\n");
printf("===============================================================================\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |b - A*x|\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];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_work, double, lwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
if ( opts.version == 2 ) {
// LAPACK complaint arguments
magma_dgeqrf2_gpu( M, N, d_A, ldda, tau, &info );
}
else {
nb = magma_get_dgeqrf_nb( M );
size = (2*min(M, N) + (N+31)/32*32 )*nb;
TESTING_MALLOC_DEV( dT, double, size );
if ( opts.version == 1 ) {
// stores dT, V blocks have zeros, R blocks inverted & stored in dT
magma_dgeqrf_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// stores dT, V blocks have zeros, R blocks stored in dT
magma_dgeqrf3_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
exit(1);
}
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check && opts.version == 2 ) {
/* =====================================================================
Check the result, following zqrt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
Only for version 2, which has LAPACK complaint output.
=================================================================== */
magma_dgetmatrix( M, N, d_A, ldda, h_R, lda );
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
double *Q, *R;
double *work;
TESTING_MALLOC_CPU( Q, double, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, double, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, double, min_mn );
// generate M by K matrix Q, where K = min(M,N)
lapackf77_dlacpy( "Lower", &M, &min_mn, h_R, &lda, Q, &ldq );
lapackf77_dorgqr( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy K by N matrix R
lapackf77_dlaset( "Lower", &min_mn, &N, &c_zero, &c_zero, R, &ldr );
lapackf77_dlacpy( "Upper", &min_mn, &N, h_R, &lda, R, &ldr );
// error = || R - Q^H*A || / (N * ||A||)
blasf77_dgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, R, &ldr );
Anorm = lapackf77_dlange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_dlange( "1", &min_mn, &N, R, &ldr, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set R = I (K by K identity), then R = I - Q^H*Q
// error = || I - Q^H*Q || / N
lapackf77_dlaset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, R, &ldr );
blasf77_dsyrk( "Upper", "Conj", &min_mn, &M, &d_neg_one, Q, &ldq, &d_one, R, &ldr );
error2 = lapackf77_dlansy( "1", "Upper", &min_mn, R, &ldr, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
else if ( opts.check && M >= N ) {
/* =====================================================================
Check the result by solving consistent linear system, A*x = b.
Only for versions 1 & 3 with M >= N.
=================================================================== */
magma_int_t lwork;
double *x, *b, *hwork;
magmaDouble_ptr d_B;
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const magma_int_t ione = 1;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, double, N );
TESTING_MALLOC_CPU( b, double, M );
lapackf77_dlarnv( &ione, ISEED, &N, x );
blasf77_dgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, double, M );
magma_dsetvector( M, b, 1, d_B, 1 );
if ( opts.version == 1 ) {
// allocate hwork
magma_dgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, double, lwork );
// solve linear system
magma_dgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork, &info );
if (info != 0)
printf("magma_dgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_dgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, double, lwork );
// solve linear system
magma_dgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork, &info );
if (info != 0)
printf("magma_dgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
exit(1);
}
magma_dgetvector( N, d_B, 1, x, 1 );
// compute r = Ax - b, saved in b
blasf77_dgemv( "Notrans", &M, &N, &c_one, h_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_dlange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_dlange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_dlange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (N * norm_A * norm_x);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check ) {
if ( opts.version == 2 ) {
bool okay = (error < tol && error2 < tol);
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else if ( M >= N ) {
bool okay = (error < tol);
status += ! okay;
printf( "%10.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf( "(error check only for M >= N)\n" );
}
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version != 2 )
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
DLATRD reduces NB rows and columns of a real symmetric matrix A to
symmetric 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, DLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, DLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by DSYTRD.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
symmetric 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 DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric 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 DOUBLE_PRECISION 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 DOUBLE_PRECISION 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 DOUBLE_PRECISION 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 real scalar, and v is a real 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 real scalar, and v is a real 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 symmetric 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_dsyev_aux
********************************************************************/
extern "C" double
magma_dlatrd_mgpu(magma_int_t num_gpus, magma_uplo_t uplo,
magma_int_t n0, magma_int_t n, magma_int_t nb, magma_int_t nb0,
double *A, magma_int_t lda,
double *e, double *tau,
double *W, magma_int_t ldw,
double **dA, magma_int_t ldda, magma_int_t offset,
double **dW, magma_int_t lddw,
double *dwork[MagmaMaxGPUs], magma_int_t ldwork,
magma_int_t k,
double *dx[MagmaMaxGPUs],
double *dy[MagmaMaxGPUs],
double *work,
magma_queue_t stream[][10],
double *times)
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(id, i, j) (dA[(id)] + ((j)+loffset)*ldda + (i) + offset)
#define dW(id, i, j) (dW[(id)] + (j) *lddw + (i))
#define dW1(id, i, j) (dW[(id)] + ((j)+nb) *lddw + (i))
double mv_time = 0.0;
magma_int_t i;
#ifndef MAGMABLAS_DSYMV_MGPU
magma_int_t loffset = nb0*((offset/nb0)/num_gpus);
#endif
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
double value = MAGMA_D_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;
double alpha;
double *dx2[MagmaMaxGPUs];
double *f;
magma_dmalloc_cpu( &f, n );
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 (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) */
double wii = *W(i, iw+1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_daxpy(&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_dlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_daxpy(&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_dlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_D_REAL( alpha );
*A(i-1,i) = MAGMA_D_MAKE( 1, 0 );
for( id=0; id < num_gpus; id++ ) {
magma_setdevice(id);
dx2[id] = dW1(id, 0, iw);
magma_dsetvector_async( n, A(0,i), 1, dW1(id, 0, iw), 1, stream[id][0]);
#ifndef MAGMABLAS_DSYMV_MGPU
magma_dsetvector_async( i, A(0,i), 1, dx[id], 1, stream[id][0] );
#endif
}
magmablas_dsymv_mgpu(num_gpus, k, MagmaUpper, 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_dgemv(MagmaTransStr, &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_dlacgv(&im1_n, W(im1, iw+1), &ldw);
#endif
blasf77_dgemv("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_dlacgv(&im1_n, W(im1, iw+1), &ldw);
lapackf77_dlacgv(&im1_n, A(im1, i +1), &lda);
#endif
blasf77_dgemv("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_dlacgv(&im1_n, A(im1, i+1), &lda);
#endif
}
// 3. Here is where we need it // TODO find the right place
magmablas_dsymv_sync(num_gpus, k, i, work, W(0, iw), stream );
if (i < n-1) {
blasf77_dgemv("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_dgemv(MagmaTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_dgemv("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_dscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_ddot_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_ddot( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -.5f * value;
blasf77_daxpy(&i, &alpha, A(0, i), &ione, W(0, iw), &ione);
for( id=0; id < num_gpus; id++ ) {
magma_setdevice(id);
if ( k > 1 ) {
magma_dsetvector_async( n, W(0,iw), 1, dW(id, 0, iw), 1, stream[id][1] );
} else {
magma_dsetvector_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" );
double wii = *W(i, i-1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_dlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_daxpy( &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_dlacgv(&i_one, &wii, &lda);
#endif
wii = -wii;
blasf77_daxpy( &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_dlarfg(&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_D_REAL( alpha );
*A(i+1,i) = MAGMA_D_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_DSYMV_MGPU
magma_setdevice(idw);
magma_dsetvector( 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_DSYMV_MGPU
dx2[id] = dW1(id, 0, i)-offset;
#else
dx2[id] = dx[id];
magma_dsetvector( i_n, A(i+1,i), 1, dx[id], 1 );
#endif
magma_dsetvector_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_dsymv_mgpu(num_gpus, k, MagmaLower, 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_dgemv(MagmaTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_dgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_dgemv(MagmaTransStr, &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_dlacgv(&i, W(ip1, 0), &ldw);
#endif
blasf77_dgemv("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_dlacgv(&i, W(ip1, 0), &ldw);
lapackf77_dlacgv(&i, A(ip1, 0), &lda);
#endif
blasf77_dgemv("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_dlacgv(&i, A(ip1, 0), &lda);
#endif
trace_cpu_end( 0 );
}
/* synchronize */
magmablas_dsymv_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_daxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_dgemv("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_dscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_ddot_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_ddot( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i]* -.5f * value;
blasf77_daxpy(&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_dsetvector_async( n, W(0,i), 1, dW(id, 0, i), 1, stream[id][1] );
} else {
magma_dsetvector_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]);
}
magma_free_cpu(f);
return mv_time;
} /* magma_dlatrd_mgpu */
extern "C" magma_int_t
magma_dlahr2(magma_int_t n, magma_int_t k, magma_int_t nb,
double *da, double *dv,
double *a, magma_int_t lda,
double *tau, double *t, magma_int_t ldt,
double *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
=======
DLAHR2 reduces the first NB columns of a real 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 DGEHRD.
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) DOUBLE_PRECISION 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) DOUBLE_PRECISION 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) DOUBLE_PRECISION array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) DOUBLE_PRECISION array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) DOUBLE_PRECISION 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 real scalar, and v is a real 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.
===================================================================== */
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
double c_neg_one = MAGMA_D_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;
double d__1;
magma_int_t i__;
double 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_dlacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
blasf77_dcopy(&i__3, &a[k+i__-1+a_dim1], &lda, &t[nb*t_dim1+1], &c__1);
blasf77_dtrmv("u","n","n",&i__3,&t[t_offset], &ldt, &t[nb*t_dim1+1], &c__1);
blasf77_dgemv("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_dlacgv(&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_dcopy(&i__2, &a[k+1+i__*a_dim1], &c__1, &t[nb*t_dim1+1], &c__1);
blasf77_dtrmv("Lower", MagmaTransStr, "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_dgemv(MagmaTransStr, &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_dtrmv("U", MagmaTransStr, "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_dgemv("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_dtrmv("L","N","U",&i__2,&a[k+1+a_dim1],&lda,&t[nb*t_dim1+1],&c__1);
blasf77_daxpy(&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_dlarfg(&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_dsetvector( i__3,
&a[k + i__ + i__*a_dim1], 1,
dv+(i__-1)*(ldda+1), 1 );
magma_dgemv(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_dgemv(MagmaTransStr, &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_D_NEGATE( tau[i__] );
blasf77_dscal(&i__2, &d__1, &t[i__ * t_dim1 + 1], &c__1);
blasf77_dtrmv("U","N","N", &i__2, &t[t_offset], &ldt, &t[i__*t_dim1+1], &c__1);
t[i__ + i__ * t_dim1] = tau[i__];
magma_dgetvector( 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_dlahr2 */
