/**
Purpose
-------
ZLATRD reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in]
nb0 INTEGER
The block size used for the matrix distribution.
nb and nb0 can be different for the final step of zhetrd.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX_16 array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
@param
dA
@param[in]
ldda
@param[in]
offset
@param
dW
@param[in]
lddw
@param
hwork
@param[in]
lhwork
@param
dwork
@param[in]
ldwork
@param[in]
queues magma_queue_t array of dimension (ngpu).
queues[dev] is an execution queue on GPU dev.
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo,
magma_int_t n, magma_int_t nb, magma_int_t nb0,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex_ptr dA[], magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex_ptr dW[], magma_int_t lddw,
magmaDoubleComplex *hwork, magma_int_t lhwork,
magmaDoubleComplex_ptr dwork[], magma_int_t ldwork,
magma_queue_t queues[] )
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(dev, i, j) (dA[(dev)] + ((j)+loffset)*ldda + (i) + offset)
#define dW(dev, i, j) (dW[(dev)] + (j) *lddw + (i))
#define dW1(dev, i, j) (dW[(dev)] + ((j)+nb) *lddw + (i))
/* Constants */
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magma_int_t ione = 1;
/* Local variables */
magmaDoubleComplex alpha, value;
magma_int_t dev;
magma_int_t i, n_i, n_i_1, ip1, iw;
// TODO check arguments
magma_int_t info = 0;
if (n <= 0) {
return info;
}
// TODO allocate f in zhetrd and pass into zlatrd. (e.g., expand hwork a bit)
magmaDoubleComplex *f;
magma_zmalloc_cpu( &f, n );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
return info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
ip1 = i + 1;
n_i_1 = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
magmaDoubleComplex wii = -conj( *W(i, iw+1) );
blasf77_zaxpy( &ip1, &wii, A(0, i+1), &ione, A(0, i), &ione );
wii = -conj( *A(i, i+1) );
blasf77_zaxpy( &ip1, &wii, W(0, iw+1), &ione, A(0, i), &ione );
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_ONE;
// TODO Previously, this set dx2[dev] = dW1(dev, 0, iw); and used dx2 in zhemv.
// TODO Now zhemv handles broadcasting x to the GPUs, but data in dW1 is
// TODO apparently still used in zhetrd_mgpu / zher2k_mgpu.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, A(0,i), 1, dW1(dev, 0, iw), 1, queues[dev] );
}
magmablas_zhemv_mgpu(
MagmaUpper, i, c_one, dA, ldda, 0,
A(0,i), 1, c_zero, W(0, iw), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
if (i < n-1) {
blasf77_zgemv( MagmaConjTransStr, &i, &n_i_1, &c_one,
W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero,
W(i+1, iw), &ione );
}
/* overlap update */
if ( i < n-1 && i-1 >= n - nb ) {
/* Update A(1:i,i) */
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, W(i-1, iw+1), &ldw );
#endif
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
A(0, i+1), &lda,
W(i-1, iw+1), &ldw, &c_one,
A(0, i-1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, W(i-1, iw+1), &ldw );
lapackf77_zlacgv( &n_i_1, A(i-1, i +1), &lda );
#endif
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
W(0, iw+1), &ldw,
A(i-1, i+1), &lda, &c_one,
A(0, i-1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &n_i_1, A(i-1, i+1), &lda );
#endif
}
// synchronize to get zhemv result W(0, iw)
magmablas_zhemv_mgpu_sync(
MagmaUpper, i, c_one, dA, ldda, 0,
A(0,i), 1, c_zero, W(0, iw), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
if (i < n-1) {
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one,
W(0, iw), &ione );
blasf77_zgemv( MagmaConjTransStr, &i, &n_i_1, &c_one,
A(0, i+1), &lda,
A(0, i), &ione, &c_zero,
W(i+1, iw), &ione );
blasf77_zgemv( "No transpose", &i, &n_i_1, &c_neg_one,
W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one,
W(0, iw), &ione );
}
blasf77_zscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy( &i, &alpha, A(0, i), &ione, W(0, iw), &ione );
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, W(0,iw), 1, dW(dev, 0, iw), 1, queues[dev] );
}
}
}
} else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
n_i = n - i;
//idw = ((offset+i)/nb)%ngpu;
if ( i > 0 ) {
trace_cpu_start( 0, "gemv", "gemv" );
magmaDoubleComplex wii = -conj( *W(i, i-1) );
blasf77_zaxpy( &n_i, &wii, A(i, i-1), &ione, A(i, i), &ione );
wii = -conj( *A(i, i-1) );
blasf77_zaxpy( &n_i, &wii, W(i, i-1), &ione, A(i, i), &ione );
}
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
n_i_1 = n - i - 1;
trace_cpu_start( 0, "larfg", "larfg" );
alpha = *A(i+1, i);
lapackf77_zlarfg( &n_i_1, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_ONE;
trace_cpu_end( 0 );
/* Compute W(i+1:n,i) */
// TODO Previously, this set dx2[id] = dW1(id, 0, i)-offset; and used dx2 in zhemv.
// TODO Now zhemv handles broadcasting x to the GPUs, but data in dW1 is
// TODO apparently still used in zhetrd_mgpu / zher2k_mgpu.
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, A(0,i), 1, dW1(dev, 0, i), 1, queues[dev] );
}
magmablas_zhemv_mgpu(
MagmaLower, n_i_1, c_one, dA, ldda, offset+i+1,
A(i+1, i), 1, c_zero, W(i+1, i), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
trace_cpu_start( 0, "gemv", "gemv" );
blasf77_zgemv( MagmaConjTransStr, &n_i_1, &i, &c_one,
W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero,
W(0, i), &ione );
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
A(i+1, 0), &lda,
W(0, i), &ione, &c_zero,
f, &ione );
blasf77_zgemv( MagmaConjTransStr, &n_i_1, &i, &c_one,
A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero,
W(0, i), &ione );
trace_cpu_end( 0 );
/* overlap update */
if ( i > 0 && i+1 < n ) {
trace_cpu_start( 0, "gemv", "gemv" );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i+1, 0), &ldw );
#endif
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
A(i+1, 0), &lda,
W(i+1, 0), &ldw, &c_one,
A(i+1, i+1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i+1, 0), &ldw );
lapackf77_zlacgv( &i, A(i+1, 0), &lda );
#endif
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
W(i+1, 0), &ldw,
A(i+1, 0), &lda, &c_one,
A(i+1, i+1), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, A(i+1, 0), &lda );
#endif
trace_cpu_end( 0 );
}
// synchronize to get zhemv result W(i+1, i)
magmablas_zhemv_mgpu_sync(
MagmaLower, n_i_1, c_one, dA, ldda, offset+i+1,
A(i+1, i), 1, c_zero, W(i+1, i), 1,
hwork, lhwork, dwork, ldwork, ngpu, nb0, queues );
trace_cpu_start( 0, "axpy", "axpy" );
if (i != 0) {
blasf77_zaxpy( &n_i_1, &c_one, f, &ione, W(i+1, i), &ione );
}
blasf77_zgemv( "No transpose", &n_i_1, &i, &c_neg_one,
W(i+1, 0), &ldw,
W(0, i), &ione, &c_one,
W(i+1, i), &ione );
blasf77_zscal( &n_i_1, &tau[i], W(i+1,i), &ione );
value = magma_cblas_zdotc( n_i_1, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy( &n_i_1, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
trace_cpu_end( 0 );
for( dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
magma_zsetvector_async( n, W(0,i), 1, dW(dev, 0, i), 1, queues[dev] );
}
}
}
}
magma_free_cpu( f );
magma_setdevice( orig_dev );
return info;
} /* magma_zlatrd_mgpu */
// --------------------
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_time=0, cpu_perf=0, gpu_time, gpu_perf, mgpu_time, mgpu_perf, cuda_time, cuda_perf;
double Ynorm, error=0, error2=0, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t n_local[MagmaMaxGPUs];
magma_int_t N, Noffset, lda, ldda, blocks, lhwork, ldwork, matsize, vecsize;
magma_int_t incx = 1;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Ylapack, *Ycublas, *Ymagma, *Ymagma1, *hwork;
magmaDoubleComplex_ptr dA, dX, dY;
magmaDoubleComplex_ptr d_lA[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magma_device_t dev;
magma_queue_t queues[MagmaMaxGPUs];
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code
double tol = opts.tolerance * lapackf77_dlamch("E");
magma_int_t nb = 64; // required by magmablas_zhemv_mgpu implementation
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_queue_create( dev, &queues[dev] );
}
// currently, tests all offsets in the offsets array;
// comment out loop below to test a specific offset.
magma_int_t offset = opts.offset;
magma_int_t offsets[] = { 0, 1, 31, 32, 33, 63, 64, 65, 100, 200 };
magma_int_t noffsets = sizeof(offsets) / sizeof(*offsets);
printf("%% uplo = %s, ngpu %d, block size = %d, offset %d\n",
lapack_uplo_const(opts.uplo), (int) opts.ngpu, (int) nb, (int) offset );
printf( "%% BLAS CUBLAS MAGMA 1 GPU MAGMA MGPU Error rel Error rel\n"
"%% N offset Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) Gflop/s (msec) to CUBLAS to LAPACK\n"
"%%==================================================================================================================\n" );
for( int itest = 0; itest < opts.ntest; ++itest ) {
// comment out these two lines & end of loop to test a specific offset
for( int ioffset=0; ioffset < noffsets; ioffset += 1 ) {
offset = offsets[ioffset];
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
Noffset = N + offset;
lda = Noffset;
ldda = magma_roundup( Noffset, opts.align ); // multiple of 32 by default
matsize = Noffset*ldda;
vecsize = (Noffset-1)*incx + 1;
gflops = FLOPS_ZHEMV( N ) / 1e9;
blocks = magma_ceildiv( N + (offset % nb), nb );
lhwork = N*opts.ngpu;
ldwork = ldda*(blocks + 1);
TESTING_MALLOC_CPU( A, magmaDoubleComplex, matsize );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ycublas, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ymagma1, magmaDoubleComplex, vecsize );
TESTING_MALLOC_CPU( Ylapack, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( X, magmaDoubleComplex, vecsize );
TESTING_MALLOC_PIN( hwork, magmaDoubleComplex, lhwork );
magma_setdevice( opts.device );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, matsize );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, vecsize );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, vecsize );
// TODO make magma_zmalloc_bcyclic helper function?
for( dev=0; dev < opts.ngpu; dev++ ) {
n_local[dev] = ((Noffset/nb)/opts.ngpu)*nb;
if (dev < (Noffset/nb) % opts.ngpu)
n_local[dev] += nb;
else if (dev == (Noffset/nb) % opts.ngpu)
n_local[dev] += Noffset % nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( d_lA[dev], magmaDoubleComplex, ldda*n_local[dev] );
TESTING_MALLOC_DEV( dwork[dev], magmaDoubleComplex, ldwork );
}
//////////////////////////////////////////////////////////////////////////
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &matsize, A );
magma_zmake_hermitian( Noffset, A, lda );
lapackf77_zlarnv( &ione, ISEED, &vecsize, X );
lapackf77_zlarnv( &ione, ISEED, &vecsize, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_setdevice( opts.device );
magma_zsetmatrix( Noffset, Noffset, A, lda, dA, ldda, opts.queue );
magma_zsetvector( Noffset, X, incx, dX, incx, opts.queue );
magma_zsetvector( Noffset, Y, incx, dY, incx, opts.queue );
cuda_time = magma_sync_wtime(0);
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo), N,
&alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
&beta, dY + offset, incx );
cuda_time = magma_sync_wtime(0) - cuda_time;
cuda_perf = gflops / cuda_time;
magma_zgetvector( Noffset, dY, incx, Ycublas, incx, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (1 GPU)
=================================================================== */
magma_setdevice( opts.device );
magma_zsetvector( Noffset, Y, incx, dY, incx, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
magmablas_zhemv_work( opts.uplo, N,
alpha, dA + offset + offset*ldda, ldda,
dX + offset, incx,
beta, dY + offset, incx, dwork[ opts.device ], ldwork,
opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
magma_zgetvector( Noffset, dY, incx, Ymagma1, incx, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (multi-GPU)
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( Noffset, Noffset, A, lda, d_lA, ldda, opts.ngpu, nb, queues );
blasf77_zcopy( &Noffset, Y, &incx, Ymagma, &incx );
// workspaces do NOT need to be zero -- set to NAN to prove
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magmablas_zlaset( MagmaFull, ldwork, 1, MAGMA_Z_NAN, MAGMA_Z_NAN, dwork[dev], ldwork, opts.queue );
}
lapackf77_zlaset( "Full", &lhwork, &ione, &MAGMA_Z_NAN, &MAGMA_Z_NAN, hwork, &lhwork );
mgpu_time = magma_sync_wtime(0);
magma_int_t info;
info = magmablas_zhemv_mgpu(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if (info != 0) {
printf("magmablas_zhemv_mgpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
info = magmablas_zhemv_mgpu_sync(
opts.uplo, N,
alpha,
d_lA, ldda, offset,
X + offset, incx,
beta,
Ymagma + offset, incx,
hwork, lhwork,
dwork, ldwork,
opts.ngpu, nb, queues );
if (info != 0) {
printf("magmablas_zhemv_sync returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
mgpu_time = magma_sync_wtime(0) - mgpu_time;
mgpu_perf = gflops / mgpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
blasf77_zcopy( &Noffset, Y, &incx, Ylapack, &incx );
cpu_time = magma_wtime();
blasf77_zhemv( lapack_uplo_const(opts.uplo), &N,
&alpha, A + offset + offset*lda, &lda,
X + offset, &incx,
&beta, Ylapack + offset, &incx );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Compute the Difference LAPACK vs. Magma
=================================================================== */
Ynorm = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ylapack, &incx );
error2 = lapackf77_zlange( "F", &Noffset, &ione, Ylapack, &Noffset, work ) / Ynorm;
}
/* =====================================================================
Compute the Difference Cublas vs. Magma
=================================================================== */
Ynorm = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work );
blasf77_zaxpy( &Noffset, &c_neg_one, Ymagma, &incx, Ycublas, &incx );
error = lapackf77_zlange( "F", &Noffset, &ione, Ycublas, &Noffset, work ) / Ynorm;
bool okay = (error < tol && error2 < tol);
status += ! okay;
if ( opts.lapack ) {
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) N, (int) offset,
cpu_perf, cpu_time*1000.,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, error2, (okay ? "ok" : "failed") );
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e --- %s\n",
(int) N, (int) offset,
cuda_perf, cuda_time*1000.,
gpu_perf, gpu_time*1000.,
mgpu_perf, mgpu_time*1000.,
error, (okay ? "ok" : "failed") );
}
/* Free Memory */
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_CPU( Ymagma1 );
TESTING_FREE_CPU( Ylapack );
TESTING_FREE_PIN( X );
TESTING_FREE_PIN( hwork );
magma_setdevice( opts.device );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
for( dev=0; dev < opts.ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_FREE_DEV( d_lA[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
// comment out these two lines line & top of loop test a specific offset
} // end for ioffset
printf( "\n" );
}
for( dev=0; dev < opts.ngpu; ++dev ) {
magma_queue_destroy( queues[dev] );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
extern "C" double
magma_zlatrd_mgpu(magma_int_t num_gpus, char uplo,
magma_int_t n0, magma_int_t n, magma_int_t nb, magma_int_t nb0,
magmaDoubleComplex *a, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *w, magma_int_t ldw,
magmaDoubleComplex **da, magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex **dw, magma_int_t lddw,
magmaDoubleComplex *dwork[MagmaMaxGPUs], magma_int_t ldwork,
magma_int_t k,
magmaDoubleComplex *dx[MagmaMaxGPUs], magmaDoubleComplex *dy[MagmaMaxGPUs],
magmaDoubleComplex *work,
magma_queue_t stream[][10],
double *times)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
ZLATRD reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = 'U', ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = 'L', ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the matrix A.
NB (input) INTEGER
The number of rows and columns to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
if UPLO = 'U', the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
if UPLO = 'L', the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= (1,N).
E (output) COMPLEX_16 array, dimension (N-1)
If UPLO = 'U', E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = 'L', E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
TAU (output) COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = 'U', and in TAU(1:nb) if UPLO = 'L'.
See Further Details.
W (output) COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
LDW (input) INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = 'U': if UPLO = 'L':
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
===================================================================== */
char uplo_[2] = {uplo, 0};
double mv_time = 0.0;
magma_int_t i;
#ifndef MAGMABLAS_ZHEMV_MGPU
magma_int_t loffset = nb0*((offset/nb0)/num_gpus);
#endif
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t id, idw, i_one = 1;
//magma_int_t kk;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *dx2[MagmaMaxGPUs];
magmaDoubleComplex *f = (magmaDoubleComplex *)malloc(n*sizeof(magmaDoubleComplex ));
if (n <= 0) {
return 0;
}
//#define PROFILE_SYMV
#ifdef PROFILE_SYMV
magma_event_t start, stop;
float etime;
magma_timestr_t cpu_start, cpu_end;
magma_setdevice(0);
magma_event_create( &start );
magma_event_create( &stop );
#endif
if (lapackf77_lsame(uplo_, "U")) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb ; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
magmaDoubleComplex wii = *W(i, iw+1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_zaxpy(&i_1, &wii, A(0, i+1), &i_one, A(0, i), &ione);
wii = *A(i, i+1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_zaxpy(&i_1, &wii, W(0, iw+1), &i_one, A(0, i), &ione);
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_MAKE( 1, 0 );
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
dx2[id] = dW1(id, 0, iw);
magma_zsetvector_async( n, A(0,i), 1, dW1(id, 0, iw), 1, stream[id][0]);
#ifndef MAGMABLAS_ZHEMV_MGPU
magma_zsetvector_async( i, A(0,i), 1, dx[id], 1, stream[id][0] );
#endif
}
magmablas_zhemv_mgpu(num_gpus, k, 'U', i, nb0, c_one, da, ldda, 0,
dx2, ione, c_zero, dy, ione, dwork, ldwork,
work, W(0, iw), stream );
if (i < n-1) {
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
/* overlap update */
if( i < n-1 && i-1 >= n - nb )
{
magma_int_t im1_1 = i_1 - 1;
magma_int_t im1 = i-1;
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magma_int_t im1_n = i_n + 1;
lapackf77_zlacgv(&im1_n, W(im1, iw+1), &ldw);
#endif
blasf77_zgemv("No transpose", &im1_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(im1, iw+1), &ldw, &c_one, A(0, i-1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&im1_n, W(im1, iw+1), &ldw);
lapackf77_zlacgv(&im1_n, A(im1, i +1), &lda);
#endif
blasf77_zgemv("No transpose", &im1_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(im1, i+1), &lda, &c_one, A(0, i-1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&im1_n, A(im1, i+1), &lda);
#endif
}
// 3. Here is where we need it // TODO find the right place
magmablas_zhemv_sync(num_gpus, k, i, work, W(0, iw), stream );
if (i < n-1) {
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -.5f * value;
blasf77_zaxpy(&i, &alpha, A(0, i), &ione, W(0, iw), &ione);
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
if( k > 1 ) {
magma_zsetvector_async( n, W(0,iw), 1, dW(id, 0, iw), 1, stream[id][1] );
} else {
magma_zsetvector_async( n, W(0,iw), 1, dW(id, 0, iw), 1, stream[id][0] );
}
}
}
}
} else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
idw = ((offset+i)/nb)%num_gpus;
if( i > 0 ) {
trace_cpu_start( 0, "gemv", "gemv" );
magmaDoubleComplex wii = *W(i, i-1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &ldw);
#endif
wii = -wii;
blasf77_zaxpy( &i_n, &wii, A(i, i-1), &ione, A(i, i), &ione);
wii = *A(i, i-1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_one, &wii, &lda);
#endif
wii = -wii;
blasf77_zaxpy( &i_n, &wii, W(i, i-1), &ione, A(i, i), &ione);
}
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
trace_cpu_start( 0, "larfg", "larfg" );
alpha = *A(i+1, i);
#ifdef PROFILE_SYMV
cpu_start = get_current_time();
#endif
lapackf77_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
#ifdef PROFILE_SYMV
cpu_end = get_current_time();
times[0] += GetTimerValue(cpu_start,cpu_end)/1000.0;
#endif
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_MAKE( 1, 0 );
trace_cpu_end( 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
//trace_gpu_start( idw, 0, "comm", "comm1" );
#ifndef MAGMABLAS_ZHEMV_MGPU
magma_setdevice(idw);
magma_zsetvector( i_n, A(i+1,i), 1, dA(idw, i+1, i), 1 );
#endif
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
trace_gpu_start( id, 0, "comm", "comm" );
#ifdef MAGMABLAS_ZHEMV_MGPU
dx2[id] = dW1(id, 0, i)-offset;
#else
dx2[id] = dx[id];
magma_zsetvector( i_n, A(i+1,i), 1, dx[id], 1 );
#endif
magma_zsetvector_async( n, A(0,i), 1, dW1(id, 0, i), 1, stream[id][0] );
trace_gpu_end( id, 0 );
}
/* mat-vec on multiple GPUs */
#ifdef PROFILE_SYMV
magma_setdevice(0);
magma_event_record(start, stream[0][0]);
#endif
magmablas_zhemv_mgpu(num_gpus, k, 'L', i_n, nb0, c_one, da, ldda, offset+i+1,
dx2, ione, c_zero, dy, ione, dwork, ldwork,
work, W(i+1,i), stream );
#ifdef PROFILE_SYMV
magma_setdevice(0);
magma_event_record(stop, stream[0][0]);
#endif
trace_cpu_start( 0, "gemv", "gemv" );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
trace_cpu_end( 0 );
/* overlap update */
if( i > 0 && i+1 < n )
{
magma_int_t ip1 = i+1;
trace_cpu_start( 0, "gemv", "gemv" );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(ip1, 0), &ldw);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(ip1, 0), &lda,
W(ip1, 0), &ldw, &c_one, A(ip1, ip1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(ip1, 0), &ldw);
lapackf77_zlacgv(&i, A(ip1 ,0), &lda);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(ip1, 0), &ldw,
A(ip1, 0), &lda, &c_one, A(ip1, ip1), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, A(ip1, 0), &lda);
#endif
trace_cpu_end( 0 );
}
/* synchronize */
magmablas_zhemv_sync(num_gpus, k, i_n, work, W(i+1,i), stream );
#ifdef PROFILE_SYMV
cudaEventElapsedTime(&etime, start, stop);
mv_time += (etime/1000.0);
times[1+(i_n/(n0/10))] += (etime/1000.0);
#endif
trace_cpu_start( 0, "axpy", "axpy" );
if (i!=0)
blasf77_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i]* -.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
trace_cpu_end( 0 );
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
if( k > 1 ) {
magma_zsetvector_async( n, W(0,i), 1, dW(id, 0, i), 1, stream[id][1] );
} else {
magma_zsetvector_async( n, W(0,i), 1, dW(id, 0, i), 1, stream[id][0] );
}
}
}
}
}
#ifdef PROFILE_SYMV
magma_setdevice(0);
magma_event_destory( start );
magma_event_destory( stop );
#endif
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
if( k > 1) magma_queue_sync(stream[id][1]);
}
free(f);
return mv_time;
} /* zlatrd_ */
