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
-------
ZLATRD reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX_16 array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
@param
dA TODO: dimension (ldda, n)?
@param
ldda TODO: ldda >= n?
@param
dW TODO: dimension (lddw, ??)
@param
lddw TODO: lddw >= n ??
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_zheev_aux
********************************************************************/
extern "C" magma_int_t
magma_zlatrd(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *W, magma_int_t ldw,
magmaDoubleComplex *work, magma_int_t lwork,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magmaDoubleComplex_ptr dW, magma_int_t lddw,
magma_queue_t queue )
{
#define A(i_, j_) (A + (i_) + (j_)*lda)
#define W(i_, j_) (W + (i_) + (j_)*ldw)
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#define dW(i_, j_) (dW + (i_) + (j_)*lddw)
/* Constants */
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magma_int_t ione = 1;
/* Local variables */
magmaDoubleComplex alpha, value;
magma_int_t i, i_n, i_1, iw;
/* Check arguments */
magma_int_t info = 0;
if ( uplo != MagmaLower && uplo != MagmaUpper ) {
info = -1;
} else if ( n < 0 ) {
info = -2;
} else if ( nb < 1 ) {
info = -3;
} else if ( lda < max(1,n) ) {
info = -5;
} else if ( ldw < max(1,n) ) {
info = -9;
} else if ( ldda < max(1,n) ) {
info = -11;
} else if ( lddw < max(1,n) ) {
info = -13;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Quick return if possible */
if (n == 0) {
return info;
}
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#ifdef COMPLEX
lapackf77_zlacgv( &i_n, W(i, iw+1), &ldw );
#endif
blasf77_zgemv( "No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i_n, W(i, iw+1), &ldw );
lapackf77_zlacgv( &i_n, A(i, i+1), &lda );
#endif
blasf77_zgemv( "No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i_n, A(i, i+1), &lda );
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg( &i, &alpha, A(0, i), &ione, &tau[i - 1] );
e[i-1] = MAGMA_Z_REAL( alpha );
*A(i-1,i) = MAGMA_Z_ONE;
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1, queue );
magma_zhemv( MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione, queue );
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw), ldw, queue );
if (i < n-1) {
blasf77_zgemv( MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( queue );
if (i < n-1) {
blasf77_zgemv( "No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione );
blasf77_zgemv( MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione );
blasf77_zgemv( "No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione );
}
blasf77_zscal( &i, &tau[i - 1], W(0, iw), &ione );
value = magma_cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy( &i, &alpha, A(0, i), &ione,
W(0, iw), &ione );
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i, 0), &ldw );
#endif
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, W(i, 0), &ldw );
lapackf77_zlacgv( &i, A(i, 0), &lda );
#endif
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione );
#ifdef COMPLEX
lapackf77_zlacgv( &i, A(i, 0), &lda );
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_zlarfg( &i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i] );
e[i] = MAGMA_Z_REAL( alpha );
*A(i+1,i) = MAGMA_Z_ONE;
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1, queue );
magma_zhemv( MagmaLower, i_n, c_one, dA(i+1, i+1), ldda,
dA(i+1, i), ione, c_zero, dW(i+1, i), ione, queue );
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, queue );
blasf77_zgemv( MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, work, &ione );
blasf77_zgemv( MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione );
// 3. Here is where we need it
magma_queue_sync( queue );
if (i != 0)
blasf77_zaxpy( &i_n, &c_one, work, &ione, W(i+1, i), &ione );
blasf77_zgemv( "No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione );
blasf77_zscal( &i_n, &tau[i], W(i+1,i), &ione );
value = magma_cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy( &i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione );
}
}
}
return info;
} /* magma_zlatrd */
extern "C" magma_int_t
magma_zlatrd2(char uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *a, magma_int_t lda,
double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *w, magma_int_t ldw,
magmaDoubleComplex *da, magma_int_t ldda,
magmaDoubleComplex *dw, magma_int_t lddw,
magmaDoubleComplex *dwork, magma_int_t ldwork)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZLATRD2 reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = 'U', ZLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = 'L', ZLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by ZHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
=========
UPLO (input) CHARACTER*1
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
= 'U': Upper triangular
= 'L': Lower triangular
N (input) INTEGER
The order of the matrix A.
NB (input) INTEGER
The number of rows and columns to be reduced.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
if UPLO = 'U', the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
if UPLO = 'L', the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= (1,N).
E (output) COMPLEX_16 array, dimension (N-1)
If UPLO = 'U', E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = 'L', E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
TAU (output) COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = 'U', and in TAU(1:nb) if UPLO = 'L'.
See Further Details.
W (output) COMPLEX_16 array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
LDW (input) INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = 'U': if UPLO = 'L':
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
===================================================================== */
char uplo_[2] = {uplo, 0};
magma_int_t i;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex value = MAGMA_Z_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaDoubleComplex alpha;
magmaDoubleComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_zmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside zlatrd
if (lapackf77_lsame(uplo_, "U")) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb ; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_zgemv("No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_zgemv("No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i_n, A(i, i+1), &ldw);
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_zlarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_Z_REAL( alpha );
MAGMA_Z_SET2REAL(*A(i-1, i), 1.);
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_zsetvector( i, A(0, i), 1, dA(0, i), 1 );
#if (GPUSHMEM < 200)
magma_zhemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione);
#else
magmablas_zhemv2(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork);
#endif
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( stream );
if (i < n-1) {
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_zgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_zgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_zscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_zdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_zaxpy(&i, &alpha, A(0, i), &ione,
W(0, iw), &ione);
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(i, 0), &ldw);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, W(i, 0), &ldw);
lapackf77_zlacgv(&i, A(i ,0), &lda);
#endif
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_zlacgv(&i, A(i, 0), &lda);
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_zlarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_Z_REAL( alpha );
MAGMA_Z_SET2REAL(*A(i+1, i), 1.);
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_zsetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
#if (GPUSHMEM < 200)
magma_zhemv(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione);
#else
magmablas_zhemv2('L', i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione,
dwork, ldwork);
#endif
// 2. Start putting the result back (asynchronously)
magma_zgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_zgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
// 3. Here is where we need it
magma_queue_sync( stream );
if (i!=0)
blasf77_zaxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_zgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_zscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_zdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_zdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_zaxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* zlatrd */
int main(int argc, char **argv)
{
TESTING_INIT();
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magma_int_t ione = 1;
real_Double_t atomics_perf=0, atomics_time=0;
real_Double_t gflops, magma_perf=0, magma_time=0, cublas_perf, cublas_time, cpu_perf, cpu_time;
double magma_error=0, atomics_error=0, cublas_error, work[1];
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, ldda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 1.5, -2.3 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( -0.6, 0.8 );
magmaDoubleComplex *A, *X, *Y, *Yatomics, *Ycublas, *Ymagma;
magmaDoubleComplex_ptr dA, dX, dY, dwork;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% N MAGMA Gflop/s (ms) Atomics Gflop/s CUBLAS Gflop/s CPU Gflop/s MAGMA error Atomics CUBLAS\n");
printf("%%=====================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_ZHEMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, magmaDoubleComplex, sizeA );
TESTING_MALLOC_CPU( X, magmaDoubleComplex, sizeX );
TESTING_MALLOC_CPU( Y, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Yatomics, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ycublas, magmaDoubleComplex, sizeY );
TESTING_MALLOC_CPU( Ymagma, magmaDoubleComplex, sizeY );
TESTING_MALLOC_DEV( dA, magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dX, magmaDoubleComplex, sizeX );
TESTING_MALLOC_DEV( dY, magmaDoubleComplex, sizeY );
blocks = magma_ceildiv( N, nb );
ldwork = ldda*blocks;
TESTING_MALLOC_DEV( dwork, magmaDoubleComplex, ldwork );
magmablas_zlaset( MagmaFull, ldwork, 1, MAGMA_Z_NAN, MAGMA_Z_NAN, dwork, ldwork );
magmablas_zlaset( MagmaFull, ldda, N, MAGMA_Z_NAN, MAGMA_Z_NAN, dA, ldda );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &sizeA, A );
magma_zmake_hermitian( N, A, lda );
// should not use data from the opposite triangle -- fill with NAN to check
magma_int_t N1 = N-1;
if ( opts.uplo == MagmaUpper ) {
lapackf77_zlaset( "Lower", &N1, &N1, &MAGMA_Z_NAN, &MAGMA_Z_NAN, &A[1], &lda );
}
else {
lapackf77_zlaset( "Upper", &N1, &N1, &MAGMA_Z_NAN, &MAGMA_Z_NAN, &A[lda], &lda );
}
lapackf77_zlarnv( &ione, ISEED, &sizeX, X );
lapackf77_zlarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_zsetmatrix( N, N, A, lda, dA, ldda );
magma_zsetvector( N, X, incx, dX, incx );
magma_zsetvector( N, Y, incy, dY, incy );
magmablasSetKernelStream( opts.queue ); // opts.handle also uses opts.queue
cublas_time = magma_sync_wtime( opts.queue );
#ifdef HAVE_CUBLAS
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
#else
magma_zhemv( opts.uplo, N, alpha, dA, 0, ldda, dX, 0, incx, beta, dY, 0, incy, opts.queue );
#endif
cublas_time = magma_sync_wtime( opts.queue ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetvector( N, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using CUBLAS - using atomics
=================================================================== */
#ifdef HAVE_CUBLAS
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_ALLOWED );
magma_zsetvector( N, Y, incy, dY, incy );
// sync on queue doesn't work -- need device sync or use NULL stream -- bug in CUBLAS?
atomics_time = magma_sync_wtime( NULL /*opts.queue*/ );
cublasZhemv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
atomics_time = magma_sync_wtime( NULL /*opts.queue*/ ) - atomics_time;
atomics_perf = gflops / atomics_time;
magma_zgetvector( N, dY, incy, Yatomics, incy );
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_NOT_ALLOWED );
#endif
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
#ifdef HAVE_CUBLAS
magma_zsetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( opts.queue );
if ( opts.version == 1 ) {
magmablas_zhemv_work( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy, dwork, ldwork, opts.queue );
}
else {
// non-work interface (has added overhead)
magmablas_zhemv( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy );
}
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetvector( N, dY, incy, Ymagma, incy );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_zhemv( lapack_uplo_const(opts.uplo), &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_zaxpy( &N, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_zlange( "M", &N, &ione, Ycublas, &N, work ) / N;
#ifdef HAVE_CUBLAS
blasf77_zaxpy( &N, &c_neg_one, Y, &incy, Yatomics, &incy );
atomics_error = lapackf77_zlange( "M", &N, &ione, Yatomics, &N, work ) / N;
blasf77_zaxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_zlange( "M", &N, &ione, Ymagma, &N, work ) / N;
#endif
bool okay = (magma_error < tol && cublas_error < tol && atomics_error < tol);
status += ! okay;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e %s\n",
(int) N,
magma_perf, 1000.*magma_time,
atomics_perf, 1000.*atomics_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error, atomics_error,
(okay ? "ok" : "failed"));
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Yatomics );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
magma_int_t
magmablas_zhemv_mgpu( magma_int_t num_gpus, magma_int_t k, char uplo,
magma_int_t n, magma_int_t nb,
magmaDoubleComplex alpha,
magmaDoubleComplex **da, magma_int_t ldda, magma_int_t offset,
magmaDoubleComplex **dx, magma_int_t incx,
magmaDoubleComplex beta,
magmaDoubleComplex **dy, magma_int_t incy,
magmaDoubleComplex **dwork, magma_int_t ldwork,
magmaDoubleComplex *work, magmaDoubleComplex *w,
magma_queue_t stream[][10] )
{
#define dX(id, i) (dx[(id)]+incx*(i))
#define dY(id, i, j) (dy[(id)]+incy*(i)+n*(j))
magma_int_t id;
#ifdef MAGMABLAS_ZHEMV_MGPU
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
magmablasSetKernelStream(stream[id][0]);
trace_gpu_start( id, 0, "memset", "memset" );
cudaMemset( dwork[id], 0, ldwork*sizeof(magmaDoubleComplex) );
trace_gpu_end( id, 0 );
trace_gpu_start( id, 0, "symv", "symv" );
}
if( nb == 32 ) {
magmablas_zhemv_mgpu_32_offset( uplo, offset+n, alpha, da, ldda,
dx, incx,
beta,
dy, incy,
dwork, ldwork,
num_gpus, nb, offset,
stream );
} else {
magmablas_zhemv_mgpu_offset( uplo, offset+n, alpha, da, ldda,
dx, incx,
beta,
dy, incy,
dwork, ldwork,
num_gpus, nb, offset,
stream );
}
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
trace_gpu_end( id, 0 );
magmablasSetKernelStream(NULL);
}
//magma_setdevice(0);
//magmablasSetKernelStream(stream[0][0]);
//magma_zhemv('L', n, alpha, &da[0][offset+offset*ldda], ldda, &dx[0][offset], incx, beta, &dy[0][offset], incy );
//magmablasSetKernelStream(NULL);
/* send to CPU */
magma_setdevice(0);
trace_gpu_start( 0, 0, "comm", "comm" );
magma_zgetvector_async( n, dY(0, offset, 0), 1, w, 1, stream[0][0] );
trace_gpu_end( 0, 0 );
magmablasSetKernelStream(NULL);
for( id=1; id<num_gpus; id++ ) {
magma_setdevice(id);
trace_gpu_start( id, 0, "comm", "comm" );
magma_zgetvector_async( n, dY(id, offset, 0), 1, &work[id*n], 1, stream[id][0] );
trace_gpu_end( id, 0 );
magmablasSetKernelStream(NULL);
}
#else
magmaDoubleComplex c_one = MAGMA_Z_ONE;
char uplo_[2] = {uplo, 0};
magma_int_t i, ii, j, kk, ib, ib0, i_1, i_local, idw;
magma_int_t i_0=n;
magma_int_t loffset0 = nb*(offset/(nb*num_gpus));
magma_int_t loffset1 = offset%nb;
magma_int_t loffset;
//magma_zhemv(uplo, n, alpha, da, ldda, dx, incx, beta, dy, incy );
idw = (offset/nb)%num_gpus;
for( id=0; id<num_gpus; id++ ) {
magma_setdevice(id);
magmablasSetKernelStream(stream[id][0]);
cudaMemset( dy[id], 0, n*k*sizeof(magmaDoubleComplex) );
}
if( lapackf77_lsame( uplo_, "L" ) ) {
/* the first block */
if( loffset1 > 0 ) {
id = idw;
kk = 0;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
loffset = loffset0+loffset1;
ib0 = min(nb-loffset1,n);
// diagonal
magma_zhemv(MagmaLower, ib0, c_one, dA(id, 0, 0 ), ldda,
dX(id, 0), incx, c_one, dY(id, 0, kk), incy);
// off-diagonl
if( ib0 < n ) {
for( j=ib0; j<n; j+= i_0 ) {
i_1 = min(i_0, n-j);
magma_zgemv(MagmaNoTrans, i_1, ib0, c_one, dA(id, j, 0), ldda,
dX(id, 0), incx, c_one, dY(id, j, kk), incy);
magma_zgemv(MagmaConjTrans, i_1, ib0, c_one, dA(id, j, 0), ldda,
dX(id, j), incx, c_one, dY(id, 0, kk), incy);
}
}
}
else {
ib0 = 0;
}
/* diagonal */
for( i=ib0; i<n; i+=nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = ((i+loffset1)/(nb*num_gpus))%k;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
i_local = (i+loffset1)/(nb*num_gpus);
ib = min(nb,n-i);
ii = nb*i_local;
loffset = loffset0;
if( id < idw ) loffset += nb;
magma_zhemv(MagmaLower, ib, c_one, dA(id, i, ii), ldda,
dX(id, i), incx, c_one, dY(id, i, kk), incy);
}
/* off-diagonal */
for( i=ib0; i<n-nb; i+=nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = ((i+loffset1)/(nb*num_gpus))%k;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
i_local = ((i+loffset1)/nb)/num_gpus;
ii = nb*i_local;
ib = min(nb,n-i);
loffset = loffset0;
if( id < idw ) loffset += nb;
for( j=i+ib; j<n; j+= i_0 ) {
i_1 = min(i_0, n-j);
magma_zgemv(MagmaNoTrans, i_1, ib, c_one, dA(id, j, ii), ldda,
dX(id, i), incx, c_one, dY(id, j, kk), incy);
magma_zgemv(MagmaConjTrans, i_1, ib, c_one, dA(id, j, ii), ldda,
dX(id, j), incx, c_one, dY(id, i, kk), incy);
}
}
} else { /* upper-triangular storage */
loffset = 0;
/* diagonal */
for( i=0; i<n; i+=nb ) {
id = (i/nb)%num_gpus;
kk = (i/(nb*num_gpus))%k;
ib = min(nb,n-i);
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
i_local = i/(nb*num_gpus);
ii = nb*i_local;
magma_zhemv(MagmaUpper, ib, c_one, dA(id, i, ii), ldda,
dX(id, i), incx, c_one, dY(id, i, kk), incy);
}
/* off-diagonal */
for( i=nb; i<n; i+=nb ) {
id = (i/nb)%num_gpus;
kk = (i/(nb*num_gpus))%k;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
i_local = (i/nb)/num_gpus;
ii = nb*i_local;
ib = min(nb,n-i);
magma_zgemv(MagmaNoTrans, i, ib, c_one, dA(id, 0, ii), ldda,
dX(id, i), incx, c_one, dY(id, 0, kk), incy);
magma_zgemv(MagmaConjTrans, i, ib, c_one, dA(id, 0, ii), ldda,
dX(id, 0), incx, c_one, dY(id, i, kk), incy);
}
}
/* send to CPU */
magma_setdevice(0);
magma_zgetvector_async( n, dY(0, 0, 0), 1, w, 1, stream[0][0] );
for( kk=1; kk<k; kk++ ) {
magma_zgetvector_async( n, dY(0, 0, kk), 1, &work[kk*n], 1, stream[0][kk] );
}
magmablasSetKernelStream(NULL);
for( id=1; id<num_gpus; id++ ) {
magma_setdevice(id);
for( kk=0; kk<k; kk++ ) {
magma_zgetvector_async( n, dY(id, 0, kk), 1, &work[id*k*n + kk*n], 1, stream[id][kk] );
}
magmablasSetKernelStream(NULL);
}
#endif
return 0;
}