extern "C" magma_int_t
magma_zhegst(magma_int_t itype, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *b, magma_int_t ldb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZHEGST reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U**H or L**H*A*L.
B must have been previously factorized as U**H*U or L*L**H by ZPOTRF.
Arguments
=========
ITYPE (input) INTEGER
= 1: compute inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H);
= 2 or 3: compute U*A*U**H or L**H*A*L.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored and B is factored as
U**H*U;
= 'L': Lower triangle of A is stored and B is factored as
L*L**H.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
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 INFO = 0, the transformed matrix, stored in the
same format as A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input) COMPLEX_16 array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
=====================================================================*/
char uplo_[2] = {uplo, 0};
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_half = MAGMA_Z_HALF;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex *dw;
magma_int_t ldda = n;
magma_int_t lddb = n;
double d_one = 1.0;
int upper = lapackf77_lsame(uplo_, "U");
/* Test the input parameters. */
*info = 0;
if (itype<1 || itype>3){
*info = -1;
}else if ((! upper) && (! lapackf77_lsame(uplo_, "L"))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
}else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
if (MAGMA_SUCCESS != magma_zmalloc( &dw, 2*n*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
nb = magma_get_zhegst_nb(n);
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
magma_zsetmatrix( n, n, B(0, 0), ldb, dB(0, 0), lddb );
/* Use hybrid blocked code */
if (itype==1) {
if (upper) {
/* Compute inv(U')*A*inv(U) */
for(k = 0; k<n; k+=nb){
kb = min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_zher2k(MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_ztrsm(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one ,dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
/* Compute inv(L)*A*inv(L') */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_zher2k(MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda);
}
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one ,dB(0,0), lddb,
dA(0,k), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ztrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k, k), &lda, B(k, k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
else {
/* Compute L'*A*L */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one ,dB(0,0), lddb,
dA(k,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ztrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_zgetmatrix( n, n, dA(0, 0), ldda, A(0, 0), lda );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dw );
return *info;
} /* magma_zhegst_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_zhemm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex calpha = MAGMA_Z_MAKE( 3.456, 5.678 );
magmaDoubleComplex cbeta = MAGMA_Z_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double error=0., errorbis=0., work[1];
magmaDoubleComplex *hA, *hX, *hB, *hR;
magmaDoubleComplex_ptr dA[MagmaMaxGPUs], dX[MagmaMaxGPUs], dB[MagmaMaxGPUs], dwork[MagmaMaxGPUs], hwork[MagmaMaxGPUs+1];
magmaDoubleComplex_ptr dA2;
magma_int_t M, N, size, lda, ldda, msize, nb, nstream;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
nstream = (opts.nstream > 0 ? opts.nstream : 2);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx = 0;
magma_buildconnection_mgpu(gnode, &nbcmplx, opts.ngpu);
printf("Initializing communication pattern... GPU-ncmplx %d\n\n", (int) nbcmplx);
for (int i=0; i < nbcmplx; ++i) {
int myngpu = gnode[i][MagmaMaxGPUs];
printf("cmplx %d has %d gpu ", i, myngpu);
for(int j=0; j < myngpu; ++j)
printf(" %d", (int) gnode[i][j]);
printf("\n");
}
magma_int_t nbevents = 2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t redevents[MagmaMaxGPUs][20];
magma_event_t redevents2[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
for( int d = 0; d < opts.ngpu; ++d ) {
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[d][i], cudaEventDisableTiming);
cudaEventCreateWithFlags(&redevents2[d][i], cudaEventDisableTiming);
}
}
printf( "nb %d, ngpu %d, nstream %d version %d\n", (int) nb, (int) opts.ngpu, (int) nstream, (int) opts.version );
printf(" M N nb offset CPU GFlop/s (sec) GPU GFlop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||X||\n");
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M;
ldda = ((M + 31)/32)*32;
size = lda*M;
gflops = FLOPS_ZHEMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = ldda*N*3;
magma_int_t hworksiz = lda*N;
TESTING_MALLOC_CPU( hA, magmaDoubleComplex, lda*M );
TESTING_MALLOC_CPU( hX, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( hB, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( hR, magmaDoubleComplex, lda*N );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( d );
TESTING_MALLOC_DEV( dA[d], magmaDoubleComplex, ldda*mlocal );
TESTING_MALLOC_DEV( dX[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dB[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork[d], magmaDoubleComplex, dworksiz );
TESTING_MALLOC_PIN( hwork[d], magmaDoubleComplex, hworksiz );
}
TESTING_MALLOC_PIN( hwork[opts.ngpu], magmaDoubleComplex, lda*N );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, magmaDoubleComplex, ldda*M );
}
lapackf77_zlarnv( &ione, iseed, &size, hA );
magma_zmake_hermitian( M, hA, lda );
size = lda*N;
lapackf77_zlarnv( &ione, iseed, &size, hX );
lapackf77_zlarnv( &ione, iseed, &size, hB );
lapackf77_zlacpy( "Full", &M, &N, hB, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
//magmablasSetKernelStream( streams[ d ][ 0 ] );
magma_zsetmatrix( M, N, hX, lda, dX[d], ldda );
//if (d == 0) magma_zsetmatrix( M, N, hB, lda, dB[d], ldda ); // this is wrong coz when offset != 0 the gpu who do the beta*C may be not 0 so this should be related to stdev(starting device who own i=0 first col)
magma_zsetmatrix( M, N, hB, lda, dB[d], ldda );
}
//memset(hR, 0, lda*N*sizeof(magmaDoubleComplex));
trace_init( 1, opts.ngpu, nstream, (magma_queue_t*) streams );
//magma_int_t offset = 0; //nb;
gpu_time = magma_sync_wtime(0);
magmablas_zhemm_mgpu_com(
MagmaLeft, MagmaLower, msize, N,
calpha, dA, ldda, offset,
dX, ldda,
cbeta, dB, ldda, dwork, dworksiz, hR, lda, hwork, hworksiz,
opts.ngpu, nb, streams, nstream, redevents2, nbevents, gnode, nbcmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "zhemm-m%d-n%d-nb%d-stream%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) nstream, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magmablasSetKernelStream( 0 );
magma_zsetmatrix( M, M, hA, lda, dA2, ldda );
magma_zsetmatrix( M, N, hX, lda, dX[0], ldda );
magma_zsetmatrix( M, N, hB, lda, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0);
magma_zhemm(
MagmaLeft, MagmaLower, msize, N,
calpha, dA2+offset*ldda+offset, ldda,
dX[0], ldda,
cbeta, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||X||
errorbis = lapackf77_zlange("fro", &msize, &msize, hA+offset*lda+offset, &lda, work );
errorbis *= lapackf77_zlange("fro", &msize, &N, hX, &lda, work );
//printf( "A =" ); magma_zprint( M, M, hA, lda );
//printf( "X =" ); magma_zprint( M, N, hX, lda );
//printf( "B =" ); magma_zprint( M, N, hB, lda );
cpu_time = magma_wtime();
blasf77_zhemm( "Left", "Lower", &msize, &N,
&calpha, hA+offset*lda+offset, &lda,
hX, &lda,
&cbeta, hB, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/*
trace_file = fopen("AJETE/C", "w");
for (int j = 0; j < N; j++)
for (int i = 0; i < siz; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, hB[j*lda+i]);
fclose(trace_file);
*/
magma_int_t firstprint=0;
for(magma_int_t dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_zgetmatrix( M, N, dB[dev], ldda, hR, lda );
// compute relative error ||R||/||A||*||X||, where R := B_magma - B_lapack = R - B
size = lda*N;
blasf77_zaxpy( &size, &c_neg_one, hB, &ione, hR, &ione );
error = lapackf77_zlange("fro", &msize, &N, hR, &lda, work) / errorbis;
//printf( "R =" ); magma_zprint( M, N, hR, lda );
if (firstprint == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (error < tol ? "ok" : "failed") );
}
else {
printf( "%89s %8.2e %s\n", " ",
error, (error < tol ? "ok" : "failed") );
}
status += ! (error < tol);
firstprint =1;
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hX );
TESTING_FREE_CPU( hB );
TESTING_FREE_PIN( hR );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
TESTING_FREE_DEV( dA[d] );
TESTING_FREE_DEV( dX[d] );
TESTING_FREE_DEV( dB[d] );
TESTING_FREE_DEV( dwork[d] );
TESTING_FREE_PIN( hwork[d] );
}
TESTING_FREE_PIN( hwork[opts.ngpu] );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( redevents[d][i] );
magma_event_destroy( redevents2[d][i] );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_zcposv_gpu(char uplo, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
magmaDoubleComplex *dworkd, magmaFloatComplex *dworks,
magma_int_t *iter, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZCPOSV computes the solution to a complex system of linear equations
A * X = B,
where A is an N-by-N Hermitian positive definite matrix and X and B
are N-by-NRHS matrices.
ZCPOSV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
=========
UPLO (input) CHARACTER
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
dA (input or input/output) COMPLEX_16 array on the GPU, dimension (LDDA,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 iterative refinement has been successfully used
(INFO.EQ.0 and ITER.GE.0, see description below), then A is
unchanged, if double factorization has been used
(INFO.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factor U or L from the Cholesky
factorization A = U**T*U or A = L*L**T.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
dB (input) COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
The N-by-NRHS right hand side matrix B.
LDDB (input) INTEGER
The leading dimension of the array dB. LDDB >= max(1,N).
dX (output) COMPLEX_16 array on the GPU, dimension (LDDX,NRHS)
If INFO = 0, the N-by-NRHS solution matrix X.
LDDX (input) INTEGER
The leading dimension of the array dX. LDDX >= max(1,N).
dworkd (workspace) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
dworks (workspace) COMPLEX array on the GPU, dimension (N*(N+NRHS))
This array is used to store the complex single precision matrix
and the right-hand sides or solutions in single precision.
ITER (output) INTEGER
< 0: iterative refinement has failed, double precision
factorization has been performed
-1 : the routine fell back to full precision for
implementation- or machine-specific reasons
-2 : narrowing the precision induced an overflow,
the routine fell back to full precision
-3 : failure of SPOTRF
-31: stop the iterative refinement after the 30th iteration
> 0: iterative refinement has been successfully used.
Returns the number of iterations
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, the leading minor of order i of (DOUBLE
PRECISION) A is not positive definite, so the
factorization could not be completed, and the solution
has not been computed.
===================================================================== */
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
#define dSX(i,j) (dSX + (i) + (j)*lddsx)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dR;
magmaFloatComplex *dSA, *dSX;
magmaDoubleComplex Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddsx = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlanhe('I', uplo, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( n, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlat2c( uplo, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cpotrf_gpu( uplo, n, dSA, lddsa, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// residual dR = dB - dA*dX in double precision
magmablas_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_zlag2c( n, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR = dB
for( j=0; j < nrhs; j++ ) {
magmablas_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER>0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_zpotrf_gpu( uplo, n, dA, ldda, info );
if (*info == 0) {
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_zpotrs_gpu( uplo, n, nrhs, dA, ldda, dX, lddx, info );
}
return *info;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaDoubleComplex *A, *B, *C, *C2, *LU;
magmaDoubleComplex *dA, *dB, *dC1, *dC2;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.5, 0.1 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_int_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
printf("=========================================================================\n");
printf( "m=%d, n=%d, k=%d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_zmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_zmalloc( &dA, size ); assert( err == 0 );
err = magma_zmalloc( &dB, size ); assert( err == 0 );
err = magma_zmalloc( &dC1, size ); assert( err == 0 );
err = magma_zmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_zlarnv( &ione, ISEED, &size, A );
lapackf77_zlarnv( &ione, ISEED, &size, B );
lapackf77_zlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test ZSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetmatrix( m, n, A, ld, dB, ld );
magma_zswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_zswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_zgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "zswap diff %.2g\n", error );
}
else {
printf( "zswap skipped for n < 3\n" );
}
// ----- test IZAMAX
// get argmax of column of A
magma_zsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_izamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIzamax( handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "izamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test ZGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetvector( maxn, B, 1, dB, 1 );
magma_zsetvector( maxn, C, 1, dC1, 1 );
magma_zsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemv( handle, cublas_trans_const(trans[ia]),
m, n, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == MagmaNoTrans ? m : n);
cublasZaxpy( handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMV( m, n ) / 1e9;
printf( "zgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZHEMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetvector( m, B, 1, dB, 1 );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zhemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemv( handle, cublas_uplo_const(uplo[iu]),
m, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMV( m ) / 1e9;
printf( "zhemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_zlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_zgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_zsetmatrix( m, m, LU, ld, dA, ld );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ztrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsv( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ztrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test ZGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_zsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_zsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemm( handle, cublas_trans_const(trans[ia]), cublas_trans_const(trans[ib]),
m, n, k, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMM( m, n, k ) / 1e9;
printf( "zgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHEMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetmatrix( m, n, B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zhemm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
m, n, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMM( side[is], m, n ) / 1e9;
printf( "zhemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHERK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_zsetmatrix( n, k, A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZherk( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &dalpha, dA, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHERK( k, n ) / 1e9;
printf( "zherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHER2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == MagmaNoTrans);
magma_zsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zher2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZher2k( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &alpha, dA, ld, dB, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHER2K( k, n ) / 1e9;
printf( "zher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasZtrmm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRMM( side[is], m, n ) / 1e9;
printf( "ztrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test ZTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( side[is], m, n ) / 1e9;
printf( "ztrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
fflush( stdout );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
int status = (total_error != 0.);
return status;
}
/**
Purpose
-------
ZHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,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 Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-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 >= max(1,N).
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT COMPLEX_16 array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(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+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_zheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_zhetrd_he2hb_mgpu( magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *tau,
magmaDoubleComplex *work, magma_int_t lwork,
magmaDoubleComplex *dAmgpu[], magma_int_t ldda,
magmaDoubleComplex *dTmgpu[], magma_int_t lddt,
magma_int_t ngpu, magma_int_t distblk,
magma_queue_t streams[][20], magma_int_t nstream,
magma_int_t *info)
{
#define A(a_1,a_2) ( A + ((a_2)-1)*( lda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define dT(a_0, a_1, a_2) (dTmgpu[a_0] + ((a_2)-1)*(lddt) + (a_1)-1)
#define dA(a_0, a_1, a_2) (dAmgpu[a_0] + ((a_2)-1)*(ldda) + (a_1)-1)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
double d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0, flipV=-1;
magma_int_t iblock, idev, di;
int i;
int lwkopt;
int lquery;
assert (nstream >= 3);
assert (nstream >= (ngpu+1));
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_Z_MAKE( lwkopt, 0 );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_int_t threads = magma_get_lapack_numthreads();
magma_int_t mklth = min(threads,16);
magma_set_lapack_numthreads(mklth);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx=0;
magma_buildconnection_mgpu(gnode, &nbcmplx, ngpu);
#ifdef ENABLE_DEBUG
printf(" Initializing communication pattern.... GPU-ncmplx %d\n\n", nbcmplx);
#endif
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magmaDoubleComplex *dspace[MagmaMaxGPUs];
magmaDoubleComplex *dwork[MagmaMaxGPUs], *dworkbis[MagmaMaxGPUs];
magmaDoubleComplex *dvall[MagmaMaxGPUs], *dv[MagmaMaxGPUs], *dw[MagmaMaxGPUs];
magmaDoubleComplex *workngpu[MagmaMaxGPUs+1];
magma_event_t redevents[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
magma_int_t nbevents = MagmaMaxGPUs*MagmaMaxGPUs;
magma_int_t lddv = ldda;
magma_int_t lddw = lddv;
magma_int_t dwrk2siz = ldda*nb*(ngpu+1);
magma_int_t worksiz = n*nb;
magma_int_t devworksiz = 2*nb*lddv + nb*lddw + nb*ldda + dwrk2siz; // 2*dv(dv0+dv1) + dw + dwork +dworkbis
// local allocation and stream creation
// TODO check malloc
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_zmalloc( &dspace[dev], devworksiz );
magma_zmalloc_pinned ( &workngpu[dev], worksiz);
dvall[dev] = dspace[dev];
dw[dev] = dvall[dev] + 2*nb*lddv;
dwork[dev] = dw[dev] + nb*lddw;
dworkbis[dev] = dwork[dev] + nb*ldda;
magmablasSetKernelStream( streams[ dev ][ 0 ] );
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[dev][i],cudaEventDisableTiming);
}
}
magma_zmalloc_pinned ( &workngpu[ngpu], worksiz);
magmaDoubleComplex *worktest = NULL;
//magma_zmalloc_cpu( &worktest, n*nb ); // not used
// ======================
magmaDoubleComplex *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(magmaDoubleComplex));
if (upper) {
printf("ZHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
} else {
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb) {
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ) {
// zpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
zpanel_to_q(MagmaUpper, pn-1, A(i, i+1), lda, work);
// find the device who own the panel then send it to the CPU.
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((i-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (i-1)%distblk; // local index in parent matrix
idev = ((i-1) / distblk) % ngpu; // device with this block
//printf("Receiving panel ofsize %d %d from idev %d A(%d,%d) \n",(pm+pn), pn,idev,i-1,di);
magma_setdevice( idev );
//magma_device_sync();
magma_zgetmatrix_async( (pm+pn), pn,
dA(idev, i, di+1), ldda,
A( i, i), lda, streams[ idev ][ nstream-1 ] );
/*
magma_device_sync();
cudaMemcpy2DAsync(A(i,i), lda*sizeof(magmaDoubleComplex),
dA(idev,i,di+1), ldda*sizeof(magmaDoubleComplex),
(pm+pn)*sizeof(magmaDoubleComplex), pn,
cudaMemcpyDeviceToHost, streams[ idev ][ nstream-1 ]);
*/
//magma_setdevice( 0 );
//printf("updating zher2k on A(%d,%d) of size %d %d \n",indi_old+pn_old-1,indi_old+pn_old-1,pm_old-pn_old,pn_old);
// compute ZHER2K_MGPU
magmablas_zher2k_mgpu2(
MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old,
c_neg_one, dv, pm_old, pn_old,
dw, pm_old, pn_old,
d_one, dAmgpu, ldda, indi_old+pn_old-1,
ngpu, distblk, streams, 2 );
//magma_setdevice( 0 );
magma_setdevice( idev );
magma_queue_sync( streams[idev][ nstream-1 ] );
//magma_setdevice( 0 );
zq_to_panel(MagmaUpper, pn-1, A(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
lapackf77_zgeqrf(&pm, &pn, A(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_zlarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, A(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
zpanel_to_q(MagmaUpper, pk, A(indi, indj), lda, work);
/* Send V and T from the CPU to the GPU */
// To be able to overlap the GET with the ZHER2K
// it should be done on last stream.
// TO Avoid a BUG that is overwriting the old_V
// used atthis moment by zher2k with the new_V
// send it now, we decide to have a flipflop
// vector of Vs. if step%2=0 use V[0] else use V[nb*n]
flipV = ((i-1)/nb)%2;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dv[dev] = dvall[dev] + flipV*nb*lddv;
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
// send V
magma_zsetmatrix_async( pm, pk,
A(indi, indj), lda,
dv[dev], pm, streams[dev][nstream-1] );
// Send the triangular factor T to the GPU
magma_zsetmatrix_async( pk, pk,
hT, nb,
dT(dev, 1, i), lddt, streams[dev][nstream-1] );
}
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// dwork = V T
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ nstream-1 ] );
magma_queue_sync( streams[dev][nstream-1] );
magma_zgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dv[dev], pm,
dT(dev, 1, i), lddt,
c_zero, dwork[dev], pm);
}
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND T WERE
// RECEIVED AND VT IS COMPUTED and SYR2K is done
// ===============================================
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t s = 0; s < nstream; ++s )
magma_queue_sync( streams[dev][s] );
}
// compute ZHEMM_MGPU
// The broadcast of the result done inside this function
// should be done in stream [0] because i am assuming this
// for the GEMMs below otherwise I have to SYNC over the
// Broadcasting stream.
if (ngpu == 1) {
magmablasSetKernelStream( streams[ 0 ][ 0 ] );
magma_zhemm(MagmaLeft, uplo, pm, pk,
c_one, dAmgpu[0]+(indi-1)*ldda+(indi-1), ldda,
dwork[0], pm,
c_zero, dw[0], pm);
} else {
magmablas_zhemm_mgpu_com(
MagmaLeft, uplo, pm, pk,
c_one, dAmgpu, ldda, indi-1,
dwork, pm,
c_zero, dw, pm, dworkbis, dwrk2siz, worktest, pm, workngpu, worksiz,
ngpu, distblk, streams, nstream-1, redevents, nbevents, gnode, nbcmplx);
}
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// Here we have to wait until the broadcast of ZHEMM has been done.
// Note that the broadcast should be done on stream[0] so in a way
// we can continue here on the same stream and avoid a sync
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
// magma_queue_sync( streams[dev][0] );
magma_zgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork[dev], pm,
dw[dev], pm,
c_zero, dworkbis[dev], nb);
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
magma_zgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dv[dev], pm,
dworkbis[dev], nb,
c_one, dw[dev], pm);
}
/* restore the panel it is put here to overlap with the previous GEMM*/
zq_to_panel(MagmaUpper, pk, A(indi, indj), lda, work);
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND W ARE DONE
// ===============================================
// Synchronise to be sure that W has been computed
// because next ZHER2K use streaming and may happen
// that lunch a gemm on stream 2 while stream 0
// which compute those 2 GEMM above has not been
// computed and also used for the same reason in
// the panel update below and also for the last HER2K
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[dev][0] );
}
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb) {
/* There would be next iteration;
do lookahead - update the next panel */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ nstream-1 ] );
//magma_queue_sync( streams[idev][0] ); removed because the sync has been done in the loop above
magma_zgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dv[idev], pm,
dw[idev], pm, c_one,
dA(idev, indi, di+1), ldda);
magma_zgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dw[idev], pm,
dv[idev], pm, c_one,
dA(idev, indi, di+1), ldda);
//printf("updating next panel distblk %d idev %d on A(%d,%d) of size %d %d %d \n",distblk,idev,indi-1,di,pm,pn,pn);
}
else {
/* no look-ahead as this is last iteration */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ 0 ] );
//printf("LAST ZHER2K idev %d on A(%d,%d) of size %d \n",idev, indi-1,di,pk);
magma_zher2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dv[idev], pm,
dw[idev], pm, d_one,
dA(idev, indi, di+1), ldda);
/* Send the last block to the CPU */
zpanel_to_q(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
magma_zgetmatrix( pk, pk,
dA(idev, indi, di+1), ldda,
A(n-pk+1, n-pk+1), lda );
zq_to_panel(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for (i)
}// end of LOWER
//magma_setdevice( 0 );
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_free( dspace[dev]);
magma_free_pinned(workngpu[dev]);
for( magma_int_t e = 0; e < nbevents; ++e ) {
cudaEventDestroy(redevents[dev][e]);
}
}
magma_free_pinned(workngpu[ngpu]);
magma_free_cpu(worktest);
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
work[0] = MAGMA_Z_MAKE( lwkopt, 0 );
magma_set_lapack_numthreads(threads);
return *info;
} /* magma_zhetrd_he2hb_mgpu */
/**
Purpose
-------
ZHEGST_GPU reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U^H)*A*inv(U) or inv(L)*A*inv(L^H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U^H or L^H*A*L.
B must have been previously factorized as U^H*U or L*L^H by ZPOTRF.
Arguments
---------
@param[in]
itype INTEGER
= 1: compute inv(U^H)*A*inv(U) or inv(L)*A*inv(L^H);
= 2 or 3: compute U*A*U^H or L^H*A*L.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored and B is factored as U^H*U;
- = MagmaLower: Lower triangle of A is stored and B is factored as L*L^H.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
dA COMPLEX_16 array, on the GPU device, dimension (LDDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, 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.
\n
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[in]
dB COMPLEX_16 array, on the GPU device, dimension (LDDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_zheev_comp
********************************************************************/
extern "C" magma_int_t
magma_zhegst_gpu(
magma_int_t itype, magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_const_ptr dB, magma_int_t lddb,
magma_int_t *info)
{
#define A(i_, j_) (work + (i_) + (j_)*lda )
#define B(i_, j_) (work + (i_) + (j_)*ldb + nb*ldb)
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#define dB(i_, j_) (dB + (i_) + (j_)*lddb)
/* Constants */
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_half = MAGMA_Z_HALF;
const magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
const double d_one = 1.0;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t k, kb, kb2, nb;
magma_int_t lda;
magma_int_t ldb;
magmaDoubleComplex *work;
bool upper = (uplo == MagmaUpper);
/* Test the input parameters. */
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! upper && uplo != MagmaLower) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
nb = magma_get_zhegst_nb( n );
lda = nb;
ldb = nb;
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, 2*nb*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
/* Use hybrid blocked code */
if (itype == 1) {
if (upper) {
kb = min( n, nb );
/* Compute inv(U^H)*A*inv(U) */
magma_zgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), ldb, queues[0] );
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
kb2 = min( n-k-nb, nb );
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(0,0), &lda, B(0,0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[0] );
if (k+kb < n) {
magma_ztrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish set dA(k,k)
// Start copying next B block
magma_zgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), ldb, queues[0] );
magma_zhemm( MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k,k+kb), ldda, queues[1] );
magma_zher2k( MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda, queues[1] );
// Start copying next A block
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, queues[0] );
magma_zhemm( MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k,k+kb), ldda, queues[1] );
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda, queues[1] );
}
}
}
else {
kb = min( n, nb );
/* Compute inv(L)*A*inv(L^H) */
magma_zgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), ldb, queues[0] );
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
kb2 = min( n-k-nb, nb );
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[0] );
if (k+kb < n) {
magma_ztrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish set dA(k,k)
// Start copying next B block
magma_zgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), ldb, queues[0] );
magma_zhemm( MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda, queues[1] );
magma_zher2k( MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda, queues[1] );
// Start copying next A block
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, queues[0] );
magma_zhemm( MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb,k), ldda, queues[1] );
magma_ztrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda, queues[1] );
}
}
}
}
else { // itype == 2 or 3
if (upper) {
/* Compute U*A*U^H */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), ldb, queues[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ztrmm( MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one, dB(0,0), lddb,
dA(0,k), ldda, queues[1] );
magma_zhemm( MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0,k), ldda, queues[1] );
magma_zher2k( MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda, queues[1] );
magma_zhemm( MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0,k), ldda, queues[1] );
magma_ztrmm( MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda, queues[1] );
}
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
lapackf77_zhegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[1] );
magma_queue_sync( queues[1] ); // wait for A(0,0) before getting next panel
}
}
else {
/* Compute L^H*A*L */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, queues[0] );
magma_zgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), ldb, queues[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ztrmm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one, dB(0,0), lddb,
dA(k,0), ldda, queues[1] );
magma_zhemm( MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k,0), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_zher2k( MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda, queues[1] );
magma_zhemm( MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k,0), ldda, queues[1] );
magma_ztrmm( MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda, queues[1] );
}
magma_queue_sync( queues[0] ); // finish get dA(k,k) -> A(0,0) and dB(k,k) -> B(0,0)
lapackf77_zhegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info );
magma_zsetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, queues[1] );
magma_queue_sync( queues[1] ); // wait for A(0,0) before getting next panel
}
}
}
magma_queue_sync( queues[0] );
magma_queue_sync( queues[1] );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free_pinned( work );
return *info;
} /* magma_zhegst_gpu */
/**
Purpose
-------
ZHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. n >= 0.
@param[in]
nb INTEGER
The inner blocking. nb >= 0.
@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 Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-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 >= max(1,N).
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT COMPLEX_16 array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(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+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_zheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_zhetrd_he2hb(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex *tau,
magmaDoubleComplex *work, magma_int_t lwork,
magmaDoubleComplex_ptr dT,
magma_int_t *info)
{
#ifdef HAVE_clBLAS
#define dA(a_1,a_2) (dA, (dA_offset + ((a_2)-1)*(ldda) + (a_1)-1))
#define dT(a_1) (dT, (dT_offset + ((a_1)-1)*(lddt)))
#else
#define dA(a_1,a_2) (dA + ((a_2)-1)*(ldda) + (a_1)-1)
#define dT(a_1) (dT + ((a_1)-1)*(lddt))
#endif
#define A(a_1,a_2) ( A + ((a_2)-1)*( lda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
magma_int_t ldda = magma_roundup( n, 32 );
magma_int_t lddt = nb;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
double d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0;
magma_int_t i;
magma_int_t lwkopt;
*info = 0;
bool upper = (uplo == MagmaUpper);
bool lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
work[0] = magma_zmake_lwork( lwkopt );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaDoubleComplex *dA;
if (MAGMA_SUCCESS != magma_zmalloc( &dA, (n + 2*nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// limit to 16 threads
magma_int_t orig_threads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads( min(orig_threads,16) );
/* Use the first panel of dA as work space */
magmaDoubleComplex *dwork = dA + n*ldda;
magmaDoubleComplex *dW = dwork + nb*ldda;
#ifdef TRACING
char buf[80];
#endif
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
trace_init( 1, 1, 3, queues );
lwork -= nb*nb;
magmaDoubleComplex *hT = work + lwork;
memset( hT, 0, nb*nb*sizeof(magmaDoubleComplex));
magma_event_t Pupdate_event;
cudaEventCreateWithFlags(&Pupdate_event,cudaEventDisableTiming);
//magma_event_create(&Pupdate_event);
if (upper) {
printf("ZHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
} else {
/* Copy the matrix to the GPU */
if (1 <= n-nb) {
trace_gpu_start( 0, 0, "set", "set A" );
magma_zsetmatrix_async( (n-nb), (n-nb),
A(nb+1, nb+1), lda,
dA(nb+1, nb+1), ldda, queues[0] );
trace_gpu_end( 0, 0 );
}
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb) {
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ) {
// magma_zpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
magma_zpanel_to_q(MagmaUpper, pn-1, A(i, i+1), lda, work);
trace_gpu_start( 0, 1, "get", "get panel" );
//magma_queue_sync( queues[0] );
magma_queue_wait_event(queues[1], Pupdate_event); //, 0);
magma_zgetmatrix_async( (pm+pn), pn,
dA( i, i), ldda,
A ( i, i), lda, queues[1] );
trace_gpu_end( 0, 1 );
trace_gpu_start( 0, 2, "her2k", "her2k" );
magma_zher2k( MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old, c_neg_one,
dA(indi_old+pn_old, indj_old), ldda,
dW + pn_old, pm_old, d_one,
dA(indi_old+pn_old, indi_old+pn_old), ldda, queues[0] );
trace_gpu_end( 0, 2 );
trace_cpu_start( 0, "sync", "sync on 1" );
magma_queue_sync( queues[1] );
trace_cpu_end( 0 );
magma_zq_to_panel(MagmaUpper, pn-1, A(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
#ifdef TRACING
snprintf( buf, sizeof(buf), "panel %d", i );
#endif
trace_cpu_start( 0, "geqrf", buf );
lapackf77_zgeqrf(&pm, &pn, A(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_zlarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, A(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
magma_zpanel_to_q(MagmaUpper, pk, A(indi, indj), lda, work);
trace_cpu_end( 0 );
/* Send V from the CPU to the GPU */
trace_gpu_start( 0, 0, "set", "set V and T" );
magma_zsetmatrix_async( pm, pk,
A(indi, indj), lda,
dA(indi, indj), ldda, queues[0] );
/* Send the triangular factor T to the GPU */
magma_zsetmatrix_async( pk, pk,
hT, nb,
dT(i), lddt, queues[0] );
trace_gpu_end( 0, 0 );
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
/* dwork = V T */
trace_cpu_start( 0, "sync", "sync on 0" );
// this sync is done here to be sure that the copy has been finished
// because below we made a restore magma_zq_to_panel and this restore need
// to ensure that the copy has been finished. we did it here to allow
// overlapp of restore with next gemm and symm.
magma_queue_sync( queues[0] );
trace_cpu_end( 0 );
trace_gpu_start( 0, 2, "gemm", "work = V*T" );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dA(indi, indj), ldda,
dT(i), lddt,
c_zero, dwork, pm, queues[0] );
trace_gpu_end( 0, 2 );
/* dW = X = A*V*T. dW = A*dwork */
trace_gpu_start( 0, 2, "hemm", "X = A*work" );
magma_zhemm( MagmaLeft, uplo, pm, pk,
c_one, dA(indi, indi), ldda,
dwork, pm,
c_zero, dW, pm, queues[0] );
trace_gpu_end( 0, 2 );
/* restore the panel */
magma_zq_to_panel(MagmaUpper, pk, A(indi, indj), lda, work);
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
trace_gpu_start( 0, 2, "gemm", "work = T'*V'*X" );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork, pm,
dW, pm,
c_zero, dwork + pm*nb, nb, queues[0] );
trace_gpu_end( 0, 2 );
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
trace_gpu_start( 0, 2, "gemm", "W = X - 0.5*V*(T'*V'*X)" );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dA(indi, indj), ldda,
dwork + pm*nb, nb,
c_one, dW, pm, queues[0] );
trace_gpu_end( 0, 2 );
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb) {
/* There would be next iteration;
do lookahead - update the next panel */
trace_gpu_start( 0, 2, "gemm", "gemm 4 next panel left" );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dA(indi, indj), ldda,
dW, pm, c_one,
dA(indi, indi), ldda, queues[0] );
trace_gpu_end( 0, 2 );
trace_gpu_start( 0, 2, "gemm", "gemm 5 next panel right" );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dW, pm,
dA(indi, indj), ldda, c_one,
dA(indi, indi), ldda, queues[0] );
trace_gpu_end( 0, 2 );
magma_event_record(Pupdate_event, queues[0]);
}
else {
/* no look-ahead as this is last iteration */
trace_gpu_start( 0, 2, "her2k", "her2k last iteration" );
magma_zher2k( MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dA(indi, indj), ldda,
dW, pm, d_one,
dA(indi, indi), ldda, queues[0] );
trace_gpu_end( 0, 2 );
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for (i)
/* Send the last block to the CPU */
pk = min(pm,pn);
if (1 <= n-nb) {
magma_zpanel_to_q(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
trace_gpu_start( 0, 2, "get", "get last block" );
magma_zgetmatrix( pk, pk,
dA(n-pk+1, n-pk+1), ldda,
A(n-pk+1, n-pk+1), lda, queues[0] );
trace_gpu_end( 0, 2 );
magma_zq_to_panel(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
}
}// end of LOWER
trace_finalize( "zhetrd_he2hb.svg", "trace.css" );
magma_queue_sync( queues[0] );
magma_queue_sync( queues[1] );
magma_event_destroy( Pupdate_event );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
magma_set_lapack_numthreads( orig_threads );
return *info;
} /* magma_zhetrd_he2hb */
