/***************************************************************************//**
Purpose
-------
SGERFS improves the computed solution to a system of linear equations.
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 SLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in]
dA REAL array on the GPU, dimension (ldda,N)
the N-by-N coefficient matrix A.
@param[in]
ldda INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
@param[in]
dB REAL array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
@param[in]
lddb INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
@param[in, out]
dX REAL array on the GPU, dimension (lddx,NRHS)
On entry, the solution matrix X, as computed by
SGETRS_NOPIV. On exit, the improved solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
@param
dworkd (workspace) REAL array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dAF REAL array on the GPU, dimension (ldda,n)
The factors L and U from the factorization A = L*U
as computed by SGETRF_NOPIV.
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, real
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 SGETRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
- > 0: if info = i, U(i,i) computed in REAL is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
@ingroup magma_gerfs_nopiv
*******************************************************************************/
extern "C" magma_int_t
magma_sgerfs_nopiv_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaFloat_ptr dA, magma_int_t ldda,
magmaFloat_ptr dB, magma_int_t lddb,
magmaFloat_ptr dX, magma_int_t lddx,
magmaFloat_ptr dworkd, magmaFloat_ptr dAF,
magma_int_t *iter,
magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
/* Constants */
const float c_neg_one = MAGMA_S_NEG_ONE;
const float c_one = MAGMA_S_ONE;
const magma_int_t ione = 1;
/* Local variables */
magmaFloat_ptr dR;
float Xnrmv, Rnrmv;
float Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, 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 = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
magma_queue_t queue = NULL;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
lddsa = n;
lddr = n;
dR = dworkd;
eps = lapackf77_slamch("Epsilon");
Anrm = magmablas_slange( MagmaInfNorm, n, n, dA, ldda, (magmaFloat_ptr)dworkd, n*nrhs, queue );
cte = Anrm * eps * magma_ssqrt( (float) n ) * BWDMAX;
// residual dR = dB - dA*dX in real
magmablas_slacpy( MagmaFull, n, nrhs, dB, lddb, dR, lddr, queue );
if ( nrhs == 1 ) {
magma_sgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1, queue );
}
else {
magma_sgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr, queue );
}
// TODO: use MAGMA_S_ABS( dX(i,j) ) instead of slange?
for( j=0; j < nrhs; j++ ) {
i = magma_isamax( n, dX(0,j), 1, queue ) - 1;
magma_sgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1, queue );
Xnrm = lapackf77_slange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_isamax( n, dR(0,j), 1, queue ) - 1;
magma_sgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1, queue );
Rnrm = lapackf77_slange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
//printf("Rnrm : %e, Xnrm*cte : %e\n", Rnrm, Xnrm*cte);
if ( Rnrm > Xnrm*cte ) {
goto refinement;
}
}
*iter = 0;
goto cleanup;
refinement:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// solve dAF*dX = dR
// it's okay that dR is used for both dB input and dX output.
magma_sgetrs_nopiv_gpu( trans, n, nrhs, dAF, lddsa, dR, lddr, info );
if (*info != 0) {
*iter = -3;
goto fallback;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_saxpycp( n, dR(0,j), dX(0,j), dB(0,j), queue );
}
// residual dR = dB - dA*dX in real
if ( nrhs == 1 ) {
magma_sgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1, queue );
}
else {
magma_sgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr, queue );
}
/* 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_isamax( n, dX(0,j), 1, queue ) - 1;
magma_sgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1, queue );
Xnrm = lapackf77_slange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_isamax( n, dR(0,j), 1, queue ) - 1;
magma_sgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1, queue );
Rnrm = lapackf77_slange( "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;
goto cleanup;
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. */
*iter = -ITERMAX - 1;
fallback:
/* Iterative refinement failed to converge to a
* satisfactory solution. */
cleanup:
magma_queue_destroy( queue );
return *info;
}
/**
Purpose
-------
SLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by SGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
dA REAL array on the GPU, dimension (LDDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements in rows K:N of the first NB columns are
overwritten with the matrix Y.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[out]
dV REAL array on the GPU, dimension (LDDV, NB)
On exit this n-by-nb array contains the Householder vectors of the transformation.
@param[in]
lddv INTEGER
The leading dimension of the array dV. LDDV >= max(1,N).
@param[in,out]
A REAL array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T REAL array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y REAL array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
Further Details
---------------
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
@verbatim
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_sgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_slahr2(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaFloat_ptr dA, magma_int_t ldda,
magmaFloat_ptr dV, magma_int_t lddv,
float *A, magma_int_t lda,
float *tau,
float *T, magma_int_t ldt,
float *Y, magma_int_t ldy )
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define Y(i_,j_) ( Y + (i_) + (j_)*ldy)
#define T(i_,j_) ( T + (i_) + (j_)*ldt)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dV(i_,j_) (dV + (i_) + (j_)*lddv)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
magma_int_t n_k_i_1, n_k;
float scale;
magma_int_t i;
float ei = MAGMA_S_ZERO;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (k < 0 || k > n) {
info = -2;
} else if (nb < 1 || nb > n) {
info = -3;
} else if (ldda < max(1,n)) {
info = -5;
} else if (lddv < max(1,n)) {
info = -7;
} else if (lda < max(1,n)) {
info = -9;
} else if (ldt < max(1,nb)) {
info = -12;
} else if (ldy < max(1,n)) {
info = -13;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
// adjust from 1-based indexing
k -= 1;
if (n <= 1)
return info;
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Update A(k:n-1,i); Update i-th column of A - Y * T * V'
// This updates one more row than LAPACK does (row k),
// making the block above the panel an even multiple of nb.
// Use last column of T as workspace, w.
// w(0:i-1, nb-1) = VA(k+i, 0:i-1)'
blasf77_scopy( &i,
A(k+i,0), &lda,
T(0,nb-1), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
// If real, conjugate row of V.
lapackf77_slacgv(&i, T(0,nb-1), &ione);
#endif
// w = T(0:i-1, 0:i-1) * w
blasf77_strmv( "Upper", "No trans", "No trans", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// A(k:n-1, i) -= Y(k:n-1, 0:i-1) * w
blasf77_sgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_scopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_strmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_sgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_strmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_sgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_strmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_saxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_slarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// dV(i+1:n-k-1, i) = VA(k+i+1:n-1, i)
magma_ssetvector( n_k_i_1,
A(k+i+1,i), 1,
dV(i+1,i), 1 );
// Compute Y(k+1:n,i) = A vi
// dA(k:n-1, i) = dA(k:n-1, i+1:n-k-1) * dV(i+1:n-k-1, i)
magma_sgemv( MagmaNoTrans, n_k, n_k_i_1,
c_one, dA(k,i+1), ldda,
dV(i+1,i), ione,
c_zero, dA(k,i), ione );
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_S_NEGATE( tau[i]);
blasf77_sgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_strmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// Y(k:n-1, i) = dA(k:n-1, i)
magma_sgetvector( n-k,
dA(k,i), 1,
Y(k,i), 1 );
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
return info;
} /* magma_slahr2 */
magma_int_t
magma_spgmres( magma_s_sparse_matrix A, magma_s_vector b, magma_s_vector *x,
magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par ){
// prepare solver feedback
solver_par->solver = Magma_PGMRES;
solver_par->numiter = 0;
solver_par->info = 0;
// local variables
float c_zero = MAGMA_S_ZERO, c_one = MAGMA_S_ONE,
c_mone = MAGMA_S_NEG_ONE;
magma_int_t dofs = A.num_rows;
magma_int_t i, j, k, m = 0;
magma_int_t restart = min( dofs-1, solver_par->restart );
magma_int_t ldh = restart+1;
float nom, rNorm, RNorm, nom0, betanom, r0 = 0.;
// CPU workspace
//magma_setdevice(0);
float *H, *HH, *y, *h1;
magma_smalloc_pinned( &H, (ldh+1)*ldh );
magma_smalloc_pinned( &y, ldh );
magma_smalloc_pinned( &HH, ldh*ldh );
magma_smalloc_pinned( &h1, ldh );
// GPU workspace
magma_s_vector r, q, q_t, z, z_t, t;
magma_s_vinit( &t, Magma_DEV, dofs, c_zero );
magma_s_vinit( &r, Magma_DEV, dofs, c_zero );
magma_s_vinit( &q, Magma_DEV, dofs*(ldh+1), c_zero );
magma_s_vinit( &z, Magma_DEV, dofs*(ldh+1), c_zero );
magma_s_vinit( &z_t, Magma_DEV, dofs, c_zero );
q_t.memory_location = Magma_DEV;
q_t.val = NULL;
q_t.num_rows = q_t.nnz = dofs;
float *dy, *dH = NULL;
if (MAGMA_SUCCESS != magma_smalloc( &dy, ldh ))
return MAGMA_ERR_DEVICE_ALLOC;
if (MAGMA_SUCCESS != magma_smalloc( &dH, (ldh+1)*ldh ))
return MAGMA_ERR_DEVICE_ALLOC;
// GPU stream
magma_queue_t stream[2];
magma_event_t event[1];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_create( &event[0] );
magmablasSetKernelStream(stream[0]);
magma_sscal( dofs, c_zero, x->val, 1 ); // x = 0
magma_scopy( dofs, b.val, 1, r.val, 1 ); // r = b
nom0 = betanom = magma_snrm2( dofs, r.val, 1 ); // nom0= || r||
nom = nom0 * nom0;
solver_par->init_res = nom0;
H(1,0) = MAGMA_S_MAKE( nom0, 0. );
magma_ssetvector(1, &H(1,0), 1, &dH(1,0), 1);
if ( (r0 = nom0 * RTOLERANCE ) < ATOLERANCE )
r0 = solver_par->epsilon;
if ( nom < r0 )
return MAGMA_SUCCESS;
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync(); tempo1=magma_wtime();
if( solver_par->verbose > 0 ){
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ){
for(k=1; k<=restart; k++) {
magma_scopy(dofs, r.val, 1, q(k-1), 1); // q[0] = 1.0/||r||
magma_sscal(dofs, 1./H(k,k-1), q(k-1), 1); // (to be fused)
q_t.val = q(k-1);
magmablasSetKernelStream(stream[0]);
// preconditioner
// z[k] = M^(-1) q(k)
magma_s_applyprecond_left( A, q_t, &t, precond_par );
magma_s_applyprecond_right( A, t, &z_t, precond_par );
magma_scopy(dofs, z_t.val, 1, z(k-1), 1);
// r = A q[k]
magma_s_spmv( c_one, A, z_t, c_zero, r );
// if (solver_par->ortho == Magma_MGS ) {
// modified Gram-Schmidt
for (i=1; i<=k; i++) {
H(i,k) =magma_sdot(dofs, q(i-1), 1, r.val, 1);
// H(i,k) = q[i] . r
magma_saxpy(dofs,-H(i,k), q(i-1), 1, r.val, 1);
// r = r - H(i,k) q[i]
}
H(k+1,k) = MAGMA_S_MAKE( magma_snrm2(dofs, r.val, 1), 0. ); // H(k+1,k) = ||r||
/*}else if (solver_par->ortho == Magma_FUSED_CGS ) {
// fusing sgemv with snrm2 in classical Gram-Schmidt
magmablasSetKernelStream(stream[0]);
magma_scopy(dofs, r.val, 1, q(k), 1);
// dH(1:k+1,k) = q[0:k] . r
magmablas_sgemv(MagmaTrans, dofs, k+1, c_one, q(0),
dofs, r.val, 1, c_zero, &dH(1,k), 1);
// r = r - q[0:k-1] dH(1:k,k)
magmablas_sgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
// 1) dH(k+1,k) = sqrt( dH(k+1,k) - dH(1:k,k) )
magma_scopyscale( dofs, k, r.val, q(k), &dH(1,k) );
// 2) q[k] = q[k] / dH(k+1,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_sgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
// asynch copy dH(1:(k+1),k) to H(1:(k+1),k)
} else {
// classical Gram-Schmidt (default)
// > explicitly calling magmabls
magmablasSetKernelStream(stream[0]);
magmablas_sgemv(MagmaTrans, dofs, k, c_one, q(0),
dofs, r.val, 1, c_zero, &dH(1,k), 1);
// dH(1:k,k) = q[0:k-1] . r
#ifndef SNRM2SCALE
// start copying dH(1:k,k) to H(1:k,k)
magma_event_record( event[0], stream[0] );
magma_queue_wait_event( stream[1], event[0] );
magma_sgetvector_async(k, &dH(1,k), 1, &H(1,k),
1, stream[1]);
#endif
// r = r - q[0:k-1] dH(1:k,k)
magmablas_sgemv(MagmaNoTrans, dofs, k, c_mone, q(0),
dofs, &dH(1,k), 1, c_one, r.val, 1);
#ifdef SNRM2SCALE
magma_scopy(dofs, r.val, 1, q(k), 1);
// q[k] = r / H(k,k-1)
magma_snrm2scale(dofs, q(k), dofs, &dH(k+1,k) );
// dH(k+1,k) = sqrt(r . r) and r = r / dH(k+1,k)
magma_event_record( event[0], stream[0] );
// start sending dH(1:k,k) to H(1:k,k)
magma_queue_wait_event( stream[1], event[0] );
// can we keep H(k+1,k) on GPU and combine?
magma_sgetvector_async(k+1, &dH(1,k), 1, &H(1,k), 1, stream[1]);
#else
H(k+1,k) = MAGMA_S_MAKE( magma_snrm2(dofs, r.val, 1), 0. );
// H(k+1,k) = sqrt(r . r)
if( k<solver_par->restart ){
magmablasSetKernelStream(stream[0]);
magma_scopy(dofs, r.val, 1, q(k), 1);
// q[k] = 1.0/H[k][k-1] r
magma_sscal(dofs, 1./H(k+1,k), q(k), 1);
// (to be fused)
}
#endif
}*/
/* Minimization of || b-Ax || in H_k */
for (i=1; i<=k; i++) {
HH(k,i) = magma_cblas_sdot( i+1, &H(1,k), 1, &H(1,i), 1 );
}
h1[k] = H(1,k)*H(1,0);
if (k != 1){
for (i=1; i<k; i++) {
HH(k,i) = HH(k,i)/HH(i,i);//
for (m=i+1; m<=k; m++){
HH(k,m) -= HH(k,i) * HH(m,i) * HH(i,i);
}
h1[k] -= h1[i] * HH(k,i);
}
}
y[k] = h1[k]/HH(k,k);
if (k != 1)
for (i=k-1; i>=1; i--) {
y[i] = h1[i]/HH(i,i);
for (j=i+1; j<=k; j++)
y[i] -= y[j] * HH(j,i);
}
m = k;
rNorm = fabs(MAGMA_S_REAL(H(k+1,k)));
}/* Minimization done */
// compute solution approximation
magma_ssetmatrix(m, 1, y+1, m, dy, m );
magma_sgemv(MagmaNoTrans, dofs, m, c_one, z(0), dofs, dy, 1,
c_one, x->val, 1);
// compute residual
magma_s_spmv( c_mone, A, *x, c_zero, r ); // r = - A * x
magma_saxpy(dofs, c_one, b.val, 1, r.val, 1); // r = r + b
H(1,0) = MAGMA_S_MAKE( magma_snrm2(dofs, r.val, 1), 0. );
// RNorm = H[1][0] = || r ||
RNorm = MAGMA_S_REAL( H(1,0) );
betanom = fabs(RNorm);
if( solver_par->verbose > 0 ){
magma_device_sync(); tempo2=magma_wtime();
if( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( betanom < r0 ) {
break;
}
}
magma_device_sync(); tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
float residual;
magma_sresidual( A, b, *x, &residual );
solver_par->iter_res = betanom;
solver_par->final_res = residual;
if( solver_par->numiter < solver_par->maxiter){
solver_par->info = 0;
}else if( solver_par->init_res > solver_par->final_res ){
if( solver_par->verbose > 0 ){
if( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = -2;
}
else{
if( solver_par->verbose > 0 ){
if( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) betanom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = -1;
}
// free pinned memory
magma_free_pinned( H );
magma_free_pinned( y );
magma_free_pinned( HH );
magma_free_pinned( h1 );
// free GPU memory
magma_free(dy);
if (dH != NULL ) magma_free(dH);
magma_s_vfree(&t);
magma_s_vfree(&r);
magma_s_vfree(&q);
magma_s_vfree(&z);
magma_s_vfree(&z_t);
// free GPU streams and events
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_event_destroy( event[0] );
magmablasSetKernelStream(NULL);
return MAGMA_SUCCESS;
} /* magma_spgmres */
/**
@deprecated
Purpose
-------
SLAQPS computes a step of QR factorization with column pivoting
of a real M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0
@param[in]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
dA REAL array, dimension (LDDA,N), on the GPU.
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau REAL array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 REAL array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 REAL array, dimension (N)
The vector with the exact column norms.
@param[in,out]
dauxv REAL array, dimension (NB), on the GPU
Auxiliary vector.
@param[in,out]
dF REAL array, dimension (LDDF,NB), on the GPU
Matrix F' = L*Y'*A.
@param[in]
lddf INTEGER
The leading dimension of the array F. LDDF >= max(1,N).
@ingroup magma_sgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_slaqps_gpu(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloat_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, float *tau,
float *vn1, float *vn2,
magmaFloat_ptr dauxv,
magmaFloat_ptr dF, magma_int_t lddf)
{
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define dF(i, j) (dF + (i) + (j)*(lddf))
float c_zero = MAGMA_S_MAKE( 0.,0.);
float c_one = MAGMA_S_MAKE( 1.,0.);
float c_neg_one = MAGMA_S_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
float z__1;
magma_int_t k, rk;
magmaFloat_ptr dAks;
float tauk = MAGMA_S_ZERO;
magma_int_t pvt;
float tol3z;
magma_int_t itemp;
float lsticc;
magmaFloat_ptr dlsticcs;
magma_smalloc( &dlsticcs, 1+256*(n+255)/256 );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
lsticc = 0;
k = 0;
magma_smalloc( &dAks, nb );
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran/CUBLAS isamax; pvt, k are 0-based.
pvt = k + magma_isamax( n-k, &vn1[k], ione, queue ) - 1;
if (pvt != k) {
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
magmablas_sswap( m, dA(0, pvt), ione, dA(0, k), ione, queue );
magmablas_sswap( i__1, dF(pvt, 0), lddf, dF(k, 0), lddf, queue );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
magma_sswap( 2, &vn1[pvt], n+offset, &vn1[k], n+offset, queue );
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_sgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, A(offset+nb, 0), lda,
F(k, 0), ldf,
c_one, A(offset+nb, k), ione, queue );
#else
i__1 = m - rk;
i__2 = k;
magma_sgemv( MagmaNoTrans, i__1, i__2,
c_neg_one, dA(rk, 0), ldda,
dF(k, 0), lddf,
c_one, dA(rk, k), ione, queue );
#endif
}
/* Generate elementary reflector H(k). */
magma_slarfg_gpu( m-rk, dA(rk, k), dA(rk + 1, k), &tau[k], &vn1[k], &dAks[k], queue );
/* needed to avoid the race condition */
if (k == 0) magma_ssetvector( 1, &c_one, 1, dA(rk, k), 1, queue );
else magma_scopymatrix( 1, 1, dA(offset, 0), 1, dA(rk, k), 1, queue );
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1 || k > 0) magma_sgetvector( 1, &tau[k], 1, &tauk, 1, queue );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Multiply on GPU */
magma_sgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, dA( rk, k+1 ), ldda,
dA( rk, k ), 1,
c_zero, dF( k+1, k ), 1, queue );
}
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K).
F(1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K)
:= tau(K)(A(RK:M,K+1:N)' - F(1:N,1:K-1)*A(RK:M,1:K-1)') A(RK:M,K)
so, F is (updated A)*V */
if (k > 0) {
z__1 = MAGMA_S_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_sgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(offset+nb, 0), lda,
dA(offset+nb, k), ione,
c_zero, dauxv, ione, queue );
i__1 = k;
magma_sgemv( MagmaNoTrans, n-k-1, i__1,
c_one, F(k+1,0), ldf,
dauxv, ione,
c_one, F(k+1,k), ione, queue );
#else
i__1 = m - rk;
i__2 = k;
magma_sgemv( MagmaConjTrans, i__1, i__2,
z__1, dA(rk, 0), ldda,
dA(rk, k), ione,
c_zero, dauxv, ione, queue );
/* I think we only need stricly lower-triangular part :) */
magma_sgemv( MagmaNoTrans, n-k-1, i__2,
c_one, dF(k+1,0), lddf,
dauxv, ione,
c_one, dF(k+1,k), ione, queue );
#endif
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
#ifdef RIGHT_UPDATE
/* right-looking update of rows, */
magma_sgemm( MagmaNoTrans, MagmaConjTrans, nb-k, i__1, ione,
c_neg_one, dA(rk, k ), ldda,
dF(k+1, k ), lddf,
c_one, dA(rk, k+1), ldda, queue );
#else
/* left-looking update of rows, *
* since F=A'v with original A, so no right-looking */
magma_sgemm( MagmaNoTrans, MagmaConjTrans, ione, i__1, i__2,
c_neg_one, dA(rk, 0 ), ldda,
dF(k+1,0 ), lddf,
c_one, dA(rk, k+1), ldda, queue );
#endif
}
/* Update partial column norms. */
if (rk < min(m, n+offset)-1 ) {
magmablas_snrm2_row_check_adjust( n-k-1, tol3z, &vn1[k+1], &vn2[k+1],
dA(rk,k+1), ldda, dlsticcs, queue );
//magma_device_sync();
magma_sgetvector( 1, &dlsticcs[0], 1, &lsticc, 1, queue );
}
++k;
}
magma_scopymatrix( 1, k, dAks, 1, dA(offset, 0), ldda+1, queue );
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
magma_sgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), lddf,
c_one, dA(rk+1, *kb), ldda, queue );
}
/* Recomputation of difficult columns. */
if ( lsticc > 0 ) {
// printf( " -- recompute dnorms --\n" );
magmablas_snrm2_check( m-rk-1, n-*kb, dA(rk+1,*kb), ldda,
&vn1[*kb], dlsticcs, queue );
magma_scopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb, queue );
}
magma_free( dAks );
magma_free( dlsticcs );
magma_queue_destroy( queue );
return MAGMA_SUCCESS;
} /* magma_slaqps */
extern "C" magma_int_t
magma_slahr2(
magma_int_t n, magma_int_t k, magma_int_t nb,
magmaFloat_ptr da, size_t da_offset, magma_int_t ldda,
magmaFloat_ptr dv, size_t dv_offset, magma_int_t lddv,
float *a, magma_int_t lda,
float *tau,
float *t, magma_int_t ldt,
float *y, magma_int_t ldy,
magma_queue_t queue)
{
/* -- clMAGMA auxiliary routine (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
SLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
This is an auxiliary routine called by SGEHRD.
Arguments
=========
N (input) INTEGER
The order of the matrix A.
K (input) INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
NB (input) INTEGER
The number of columns to be reduced.
DA (input/output) REAL array on the GPU, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
DV (output) REAL array on the GPU, dimension (N, NB)
On exit this contains the Householder vectors of the transformation.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
T (output) REAL array, dimension (LDT,NB)
The upper triangular matrix T.
LDT (input) INTEGER
The leading dimension of the array T. LDT >= NB.
Y (output) REAL array, dimension (LDY,NB)
The n-by-nb matrix Y.
LDY (input) INTEGER
The leading dimension of the array Y. LDY >= N.
Further Details
===============
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
===================================================================== */
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t c__1 = 1;
magma_int_t a_dim1, a_offset, t_dim1, t_offset, y_dim1, y_offset, i__2, i__3;
float d__1;
magma_int_t i__;
float ei;
--tau;
a_dim1 = lda;
a_offset = 1 + a_dim1;
a -= a_offset;
t_dim1 = ldt;
t_offset = 1 + t_dim1;
t -= t_offset;
y_dim1 = ldy;
y_offset = 1 + y_dim1;
y -= y_offset;
if (n <= 1)
return MAGMA_SUCCESS;
for (i__ = 1; i__ <= nb; ++i__) {
if (i__ > 1) {
/* Update A(K+1:N,I); Update I-th column of A - Y * V' */
i__2 = n - k + 1;
i__3 = i__ - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
blasf77_scopy(&i__3, &a[k+i__-1+a_dim1], &lda, &t[nb*t_dim1+1], &c__1);
blasf77_strmv("u","n","n",&i__3,&t[t_offset], &ldt, &t[nb*t_dim1+1], &c__1);
blasf77_sgemv("NO TRANSPOSE", &i__2, &i__3, &c_neg_one, &y[k + y_dim1],
&ldy, &t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__*a_dim1],&c__1);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv(&i__3, &a[k+i__-1+a_dim1], &lda);
#endif
/* Apply I - V * T' * V' to this column (call it b) from the
left, using the last column of T as workspace
Let V = ( V1 ) and b = ( b1 ) (first I-1 rows)
( V2 ) ( b2 )
where V1 is unit lower triangular
w := V1' * b1 */
i__2 = i__ - 1;
blasf77_scopy(&i__2, &a[k+1+i__*a_dim1], &c__1, &t[nb*t_dim1+1], &c__1);
blasf77_strmv("Lower", MagmaConjTransStr, "UNIT", &i__2,
&a[k + 1 + a_dim1], &lda, &t[nb * t_dim1 + 1], &c__1);
/* w := w + V2'*b2 */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_sgemv(MagmaConjTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_one, &t[nb*t_dim1+1], &c__1);
/* w := T'*w */
i__2 = i__ - 1;
blasf77_strmv("U", MagmaConjTransStr, "N", &i__2, &t[t_offset], &ldt,
&t[nb*t_dim1+1], &c__1);
/* b2 := b2 - V2*w */
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_sgemv("N", &i__2, &i__3, &c_neg_one, &a[k + i__ + a_dim1], &lda,
&t[nb*t_dim1+1], &c__1, &c_one, &a[k+i__+i__*a_dim1], &c__1);
/* b1 := b1 - V1*w */
i__2 = i__ - 1;
blasf77_strmv("L","N","U",&i__2,&a[k+1+a_dim1],&lda,&t[nb*t_dim1+1],&c__1);
blasf77_saxpy(&i__2, &c_neg_one, &t[nb * t_dim1 + 1], &c__1,
&a[k + 1 + i__ * a_dim1], &c__1);
a[k + i__ - 1 + (i__ - 1) * a_dim1] = ei;
}
/* Generate the elementary reflector H(I) to annihilate A(K+I+1:N,I) */
i__2 = n - k - i__ + 1;
i__3 = k + i__ + 1;
lapackf77_slarfg(&i__2, &a[k + i__ + i__ * a_dim1],
&a[min(i__3,n) + i__ * a_dim1], &c__1, &tau[i__]);
ei = a[k + i__ + i__ * a_dim1];
a[k + i__ + i__ * a_dim1] = c_one;
/* Compute Y(K+1:N,I) */
i__2 = n - k;
i__3 = n - k - i__ + 1;
magma_ssetvector( i__3, &a[k + i__ + i__*a_dim1], 1, dv, dv_offset+(i__-1)*(lddv+1), 1, queue );
magma_sgemv(MagmaNoTrans, i__2+1, i__3, c_one,
da, da_offset + (-1 + k + i__ * ldda), ldda,
dv, dv_offset + (i__-1)*(lddv+1), c__1, c_zero,
da, da_offset + (-1 + k + (i__-1)*ldda), c__1, queue);
i__2 = n - k - i__ + 1;
i__3 = i__ - 1;
blasf77_sgemv(MagmaConjTransStr, &i__2, &i__3, &c_one,
&a[k + i__ + a_dim1], &lda, &a[k+i__+i__*a_dim1], &c__1,
&c_zero, &t[i__*t_dim1+1], &c__1);
/* Compute T(1:I,I) */
i__2 = i__ - 1;
d__1 = MAGMA_S_NEGATE( tau[i__] );
blasf77_sscal(&i__2, &d__1, &t[i__ * t_dim1 + 1], &c__1);
blasf77_strmv("U","N","N", &i__2, &t[t_offset], &ldt, &t[i__*t_dim1+1], &c__1);
t[i__ + i__ * t_dim1] = tau[i__];
magma_sgetvector( n - k + 1, da, da_offset+(-1+ k+(i__-1)*ldda), 1, y+ k + i__*y_dim1, 1, queue );
}
a[k + nb + nb * a_dim1] = ei;
return MAGMA_SUCCESS;
} /* magma_slahr2 */
/**
Purpose
-------
SLAQPS computes a step of QR factorization with column pivoting
of a real M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0
@param[in]
offset INTEGER
The number of rows of A that have been factorized in
previous steps.
@param[in]
nb INTEGER
The number of columns to factorize.
@param[out]
kb INTEGER
The number of columns actually factorized.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
jpvt INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
@param[out]
tau REAL array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 REAL array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 REAL array, dimension (N)
The vector with the exact column norms.
@param[in,out]
auxv REAL array, dimension (NB)
Auxiliar vector.
@param[in,out]
F REAL array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
@param[in]
ldf INTEGER
The leading dimension of the array F. LDF >= max(1,N).
@ingroup magma_sgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_slaqps(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
float *A, magma_int_t lda,
float *dA, magma_int_t ldda,
magma_int_t *jpvt, float *tau, float *vn1, float *vn2,
float *auxv,
float *F, magma_int_t ldf,
float *dF, magma_int_t lddf)
{
#define A(i, j) (A + (i) + (j)*(lda ))
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define F(i, j) (F + (i) + (j)*(ldf ))
#define dF(i, j) (dF + (i) + (j)*(lddf))
float c_zero = MAGMA_S_MAKE( 0.,0.);
float c_one = MAGMA_S_MAKE( 1.,0.);
float c_neg_one = MAGMA_S_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
float d__1;
float z__1;
magma_int_t j, k, rk;
float Akk;
magma_int_t pvt;
float temp, temp2, tol3z;
magma_int_t itemp;
magma_int_t lsticc;
magma_int_t lastrk;
lastrk = min( m, n + offset );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
magma_queue_t stream;
magma_queue_create( &stream );
lsticc = 0;
k = 0;
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran isamax; pvt, k are 0-based.
i__1 = n-k;
pvt = k + blasf77_isamax( &i__1, &vn1[k], &ione ) - 1;
if (pvt != k) {
if (pvt >= nb) {
/* 1. Start copy from GPU */
magma_sgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, stream );
}
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
blasf77_sswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
vn1[pvt] = vn1[k];
vn2[pvt] = vn2[k];
if (pvt < nb) {
/* no need of transfer if pivot is within the panel */
blasf77_sswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
/* 1. Finish copy from GPU */
magma_queue_sync( stream );
/* 2. Swap as usual on CPU */
blasf77_sswap(&m, A(0, pvt), &ione, A(0, k), &ione);
/* 3. Restore the GPU */
magma_ssetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
#if defined(PRECISION_c) || defined(PRECISION_z)
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_S_CNJG( *F(k,j) );
}
#endif
i__1 = m - rk;
i__2 = k;
blasf77_sgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );
#if defined(PRECISION_c) || defined(PRECISION_z)
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_S_CNJG( *F(k,j) );
}
#endif
}
/* Generate elementary reflector H(k). */
if (rk < m-1) {
i__1 = m - rk;
lapackf77_slarfg( &i__1, A(rk, k), A(rk + 1, k), &ione, &tau[k] );
} else {
lapackf77_slarfg( &ione, A(rk, k), A(rk, k), &ione, &tau[k] );
}
Akk = *A(rk, k);
*A(rk, k) = c_one;
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
magma_ssetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL SGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
magma_int_t i__3 = nb-k-1;
magma_int_t i__4 = i__2 - i__3;
magma_int_t i__5 = nb-k;
magma_sgemv( MagmaConjTrans, i__1 - i__5, i__2 - i__3,
tau[k], dA(rk +i__5, k+1+i__3), ldda,
dA(rk +i__5, k ), ione,
c_zero, dF(k+1+i__3, k ), ione );
magma_sgetmatrix_async( i__2-i__3, 1,
dF(k + 1 +i__3, k), i__2,
F (k + 1 +i__3, k), i__2, stream );
blasf77_sgemv( MagmaConjTransStr, &i__1, &i__3,
&tau[k], A(rk, k+1), &lda,
A(rk, k ), &ione,
&c_zero, F(k+1, k ), &ione );
magma_queue_sync( stream );
blasf77_sgemv( MagmaConjTransStr, &i__5, &i__4,
&tau[k], A(rk, k+1+i__3), &lda,
A(rk, k ), &ione,
&c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros. */
for (j = 0; j < k; ++j) {
*F(j, k) = c_zero;
}
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K). */
if (k > 0) {
i__1 = m - rk;
i__2 = k;
z__1 = MAGMA_S_NEGATE( tau[k] );
blasf77_sgemv( MagmaConjTransStr, &i__1, &i__2,
&z__1, A(rk, 0), &lda,
A(rk, k), &ione,
&c_zero, auxv, &ione );
i__1 = k;
blasf77_sgemv( MagmaNoTransStr, &n, &i__1,
&c_one, F(0,0), &ldf,
auxv, &ione,
&c_one, F(0,k), &ione );
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
blasf77_sgemm( MagmaNoTransStr, MagmaConjTransStr, &ione, &i__1, &i__2,
&c_neg_one, A(rk, 0 ), &lda,
F(k+1,0 ), &ldf,
&c_one, A(rk, k+1), &lda );
}
/* Update partial column norms. */
if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
/* NOTE: The following 4 lines follow from the analysis in
Lapack Working Note 176. */
temp = MAGMA_S_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (float) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_ssqrt(temp);
}
}
}
}
*A(rk, k) = Akk;
++k;
}
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU */
magma_ssetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );
magma_sgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), i__2,
c_one, dA(rk+1, *kb), ldda );
}
/* Recomputation of difficult columns. */
while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb)
vn1[lsticc] = magma_cblas_snrm2( i__1, A(rk+1,lsticc), ione );
else {
/* Where is the data, CPU or GPU ? */
float r1, r2;
r1 = magma_cblas_snrm2( nb-k, A(rk+1,lsticc), ione );
r2 = magma_snrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
//vn1[lsticc] = magma_snrm2(i__1, dA(rk + 1, lsticc), ione);
vn1[lsticc] = magma_ssqrt(r1*r1 + r2*r2);
}
/* NOTE: The computation of VN1( LSTICC ) relies on the fact that
SNRM2 does not fail on vectors with norm below the value of SQRT(SLAMCH('S')) */
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;
}
magma_queue_destroy( stream );
return MAGMA_SUCCESS;
} /* magma_slaqps */
/***************************************************************************//**
Purpose
-------
SLABRD reduces the first NB rows and columns of a real general
m by n matrix A to upper or lower bidiagonal form by an orthogonal
transformation Q' * A * P, and returns the matrices X and Y which
are needed to apply the transformation to the unreduced part of A.
If m >= n, A is reduced to upper bidiagonal form; if m < n, to lower
bidiagonal form.
This is an auxiliary routine called by SGEBRD.
Arguments
---------
@param[in]
m INTEGER
The number of rows in the matrix A.
@param[in]
n INTEGER
The number of columns in the matrix A.
@param[in]
nb INTEGER
The number of leading rows and columns of A to be reduced.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the m by n general matrix to be reduced.
On exit, the first NB rows and columns of the matrix are
overwritten; the rest of the array is unchanged.
If m >= n, elements on and below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors; and
elements above the diagonal in the first NB rows, with the
array TAUP, represent the orthogonal matrix P as a product
of elementary reflectors.
\n
If m < n, elements below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors, and
elements on and above the diagonal in the first NB rows,
with the array TAUP, represent the orthogonal matrix P as
a product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
dA REAL array, dimension (LDDA,N)
Copy of A on GPU.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
d REAL array, dimension (NB)
The diagonal elements of the first NB rows and columns of
the reduced matrix. D(i) = A(i,i).
@param[out]
e REAL array, dimension (NB)
The off-diagonal elements of the first NB rows and columns of
the reduced matrix.
@param[out]
tauq REAL array dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
@param[out]
taup REAL array, dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
@param[out]
X REAL array, dimension (LDX,NB)
The m-by-nb matrix X required to update the unreduced part
of A.
@param[in]
ldx INTEGER
The leading dimension of the array X. LDX >= M.
@param[out]
dX REAL array, dimension (LDDX,NB)
Copy of X on GPU.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= M.
@param[out]
Y REAL array, dimension (LDY,NB)
The n-by-nb matrix Y required to update the unreduced part
of A.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[out]
dY REAL array, dimension (LDDY,NB)
Copy of Y on GPU.
@param[in]
lddy INTEGER
The leading dimension of the array dY. LDDY >= N.
@param
work REAL array, dimension (LWORK)
Workspace.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= max( M, N ).
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
The matrices Q and P are represented as products of elementary
reflectors:
Q = H(1) H(2) . . . H(nb) and P = G(1) G(2) . . . G(nb)
Each H(i) and G(i) has the form:
H(i) = I - tauq * v * v' and G(i) = I - taup * u * u'
where tauq and taup are real scalars, and v and u are real vectors.
If m >= n, v(1:i-1) = 0, v(i) = 1, and v(i:m) is stored on exit in
A(i:m,i); u(1:i) = 0, u(i+1) = 1, and u(i+1:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
If m < n, v(1:i) = 0, v(i+1) = 1, and v(i+1:m) is stored on exit in
A(i+2:m,i); u(1:i-1) = 0, u(i) = 1, and u(i:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
The elements of the vectors v and u together form the m-by-nb matrix
V and the nb-by-n matrix U' which are needed, with X and Y, to apply
the transformation to the unreduced part of the matrix, using a block
update of the form: A := A - V*Y' - X*U'.
The contents of A on exit are illustrated by the following examples
with nb = 2:
@verbatim
m = 6 and n = 5 (m > n): m = 5 and n = 6 (m < n):
( 1 1 u1 u1 u1 ) ( 1 u1 u1 u1 u1 u1 )
( v1 1 1 u2 u2 ) ( 1 1 u2 u2 u2 u2 )
( v1 v2 a a a ) ( v1 1 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a )
@endverbatim
where a denotes an element of the original matrix which is unchanged,
vi denotes an element of the vector defining H(i), and ui an element
of the vector defining G(i).
@ingroup magma_labrd
*******************************************************************************/
extern "C" magma_int_t
magma_slabrd_gpu(
magma_int_t m, magma_int_t n, magma_int_t nb,
float *A, magma_int_t lda,
magmaFloat_ptr dA, magma_int_t ldda,
float *d, float *e, float *tauq, float *taup,
float *X, magma_int_t ldx,
magmaFloat_ptr dX, magma_int_t lddx,
float *Y, magma_int_t ldy,
magmaFloat_ptr dY, magma_int_t lddy,
float *work, magma_int_t lwork,
magma_queue_t queue )
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define X(i_,j_) ( X + (i_) + (j_)*ldx)
#define Y(i_,j_) ( Y + (i_) + (j_)*ldy)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dY(i_,j_) (dY + (i_) + (j_)*lddy)
#define dX(i_,j_) (dX + (i_) + (j_)*lddx)
/* Constants */
const float c_neg_one = MAGMA_S_NEG_ONE;
const float c_one = MAGMA_S_ONE;
const float c_zero = MAGMA_S_ZERO;
const magma_int_t ione = 1;
/* Local variables */
magma_int_t i, i1, m_i, m_i1, n_i, n_i1;
float alpha;
/* Quick return if possible */
magma_int_t info = 0;
if (m <= 0 || n <= 0) {
return info;
}
if (m >= n) {
/* Reduce to upper bidiagonal form */
for (i=0; i < nb; ++i) {
/* Update A(i:m,i) */
i1 = i + 1;
m_i = m - i;
m_i1 = m - (i+1);
n_i1 = n - (i+1);
#ifdef COMPLEX
lapackf77_slacgv( &i, Y(i,0), &ldy );
#endif
blasf77_sgemv( "No transpose", &m_i, &i, &c_neg_one,
A(i,0), &lda,
Y(i,0), &ldy, &c_one,
A(i,i), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &i, Y(i,0), &ldy );
#endif
blasf77_sgemv( "No transpose", &m_i, &i, &c_neg_one,
X(i,0), &ldx,
A(0,i), &ione, &c_one,
A(i,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+1:m,i) */
alpha = *A(i,i);
lapackf77_slarfg( &m_i, &alpha, A(min(i+1,m-1),i), &ione, &tauq[i] );
d[i] = MAGMA_S_REAL( alpha );
if (i+1 < n) {
*A(i,i) = c_one;
/* Compute Y(i+1:n,i) */
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_ssetvector( m_i,
A(i,i), 1,
dA(i,i), 1, queue );
// 2. Multiply ---------------------------------------------
magma_sgemv( MagmaConjTrans, m_i, n_i1, c_one,
dA(i,i+1), ldda,
dA(i,i), ione, c_zero,
dY(i+1,i), ione, queue );
// 3. Get the result back ----------------------------------
magma_sgetmatrix_async( n_i1, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, queue );
blasf77_sgemv( MagmaConjTransStr, &m_i, &i, &c_one,
A(i,0), &lda,
A(i,i), &ione, &c_zero,
Y(0,i), &ione );
blasf77_sgemv( "N", &n_i1, &i, &c_neg_one,
Y(i+1,0), &ldy,
Y(0,i), &ione, &c_zero,
work, &ione );
blasf77_sgemv( MagmaConjTransStr, &m_i, &i, &c_one,
X(i,0), &ldx,
A(i,i), &ione, &c_zero,
Y(0,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( queue );
if (i != 0) {
blasf77_saxpy( &n_i1, &c_one, work, &ione, Y(i+1,i), &ione );
}
blasf77_sgemv( MagmaConjTransStr, &i, &n_i1, &c_neg_one,
A(0,i+1), &lda,
Y(0,i), &ione, &c_one,
Y(i+1,i), &ione );
blasf77_sscal( &n_i1, &tauq[i], Y(i+1,i), &ione );
/* Update A(i,i+1:n) */
#ifdef COMPLEX
lapackf77_slacgv( &n_i1, A(i,i+1), &lda );
lapackf77_slacgv( &i1, A(i,0), &lda );
#endif
blasf77_sgemv( "No transpose", &n_i1, &i1, &c_neg_one,
Y(i+1,0), &ldy,
A(i,0), &lda, &c_one,
A(i,i+1), &lda );
#ifdef COMPLEX
lapackf77_slacgv( &i1, A(i,0), &lda );
lapackf77_slacgv( &i, X(i,0), &ldx );
#endif
blasf77_sgemv( MagmaConjTransStr, &i, &n_i1, &c_neg_one,
A(0,i+1), &lda,
X(i,0), &ldx, &c_one,
A(i,i+1), &lda );
#ifdef COMPLEX
lapackf77_slacgv( &i, X(i,0), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+2:n) */
alpha = *A(i,i+1);
lapackf77_slarfg( &n_i1, &alpha, A(i,min(i+2,n-1)), &lda, &taup[i] );
e[i] = MAGMA_S_REAL( alpha );
*A(i,i+1) = c_one;
/* Compute X(i+1:m,i) */
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_ssetvector( n_i1,
A(i,i+1), lda,
dA(i,i+1), ldda, queue );
// 2. Multiply ---------------------------------------------
magma_sgemv( MagmaNoTrans, m_i1, n_i1, c_one,
dA(i+1,i+1), ldda,
dA(i,i+1), ldda,
//dY(0,0), 1,
c_zero,
dX(i+1,i), ione, queue );
// 3. Get the result back ----------------------------------
magma_sgetmatrix_async( m_i1, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, queue );
blasf77_sgemv( MagmaConjTransStr, &n_i1, &i1, &c_one,
Y(i+1,0), &ldy,
A(i,i+1), &lda, &c_zero,
X(0,i), &ione );
blasf77_sgemv( "N", &m_i1, &i1, &c_neg_one,
A(i+1,0), &lda,
X(0,i), &ione, &c_zero,
work, &ione );
blasf77_sgemv( "N", &i, &n_i1, &c_one,
A(0,i+1), &lda,
A(i,i+1), &lda, &c_zero,
X(0,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( queue );
if ((i+1) != 0) {
blasf77_saxpy( &m_i1, &c_one, work, &ione, X(i+1,i), &ione );
}
blasf77_sgemv( "No transpose", &m_i1, &i, &c_neg_one,
X(i+1,0), &ldx,
X(0,i), &ione, &c_one,
X(i+1,i), &ione );
blasf77_sscal( &m_i1, &taup[i], X(i+1,i), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &n_i1, A(i,i+1), &lda );
// 4. Send the block reflector A(i+1:m,i) to the GPU after SLACGV()
magma_ssetvector( n_i1,
A(i,i+1), lda,
dA(i,i+1), ldda, queue );
#endif
}
}
}
else {
/* Reduce to lower bidiagonal form */
for (i=0; i < nb; ++i) {
/* Update A(i,i:n) */
i1 = i + 1;
m_i1 = m - (i+1);
n_i = n - i;
n_i1 = n - (i+1);
#ifdef COMPLEX
lapackf77_slacgv( &n_i, A(i,i), &lda );
lapackf77_slacgv( &i, A(i,0), &lda );
#endif
blasf77_sgemv( "No transpose", &n_i, &i, &c_neg_one,
Y(i,0), &ldy,
A(i,0), &lda, &c_one,
A(i,i), &lda );
#ifdef COMPLEX
lapackf77_slacgv( &i, A(i,0), &lda );
lapackf77_slacgv( &i, X(i,0), &ldx );
#endif
blasf77_sgemv( MagmaConjTransStr, &i, &n_i, &c_neg_one,
A(0,i), &lda,
X(i,0), &ldx, &c_one,
A(i,i), &lda );
#ifdef COMPLEX
lapackf77_slacgv( &i, X(i,0), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+1:n) */
alpha = *A(i,i);
lapackf77_slarfg( &n_i, &alpha, A(i,min(i+1,n-1)), &lda, &taup[i] );
d[i] = MAGMA_S_REAL( alpha );
if (i+1 < m) {
*A(i,i) = c_one;
/* Compute X(i+1:m,i) */
// 1. Send the block reflector A(i,i+1:n) to the GPU ------
magma_ssetvector( n_i,
A(i,i), lda,
dA(i,i), ldda, queue );
// 2. Multiply ---------------------------------------------
magma_sgemv( MagmaNoTrans, m_i1, n_i, c_one,
dA(i+1,i), ldda,
dA(i,i), ldda,
//dY(0,0), 1,
c_zero,
dX(i+1,i), ione, queue );
// 3. Get the result back ----------------------------------
magma_sgetmatrix_async( m_i1, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, queue );
blasf77_sgemv( MagmaConjTransStr, &n_i, &i, &c_one,
Y(i,0), &ldy,
A(i,i), &lda, &c_zero,
X(0,i), &ione );
blasf77_sgemv( "No transpose", &m_i1, &i, &c_neg_one,
A(i+1,0), &lda,
X(0,i), &ione, &c_zero,
work, &ione );
blasf77_sgemv( "No transpose", &i, &n_i, &c_one,
A(0,i), &lda,
A(i,i), &lda, &c_zero,
X(0,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( queue );
if (i != 0) {
blasf77_saxpy( &m_i1, &c_one, work, &ione, X(i+1,i), &ione );
}
blasf77_sgemv( "No transpose", &m_i1, &i, &c_neg_one,
X(i+1,0), &ldx,
X(0,i), &ione, &c_one,
X(i+1,i), &ione );
blasf77_sscal( &m_i1, &taup[i], X(i+1,i), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &n_i, A(i,i), &lda );
magma_ssetvector( n_i,
A(i,i), lda,
dA(i,i), ldda, queue );
#endif
/* Update A(i+1:m,i) */
#ifdef COMPLEX
lapackf77_slacgv( &i, Y(i,0), &ldy );
#endif
blasf77_sgemv( "No transpose", &m_i1, &i, &c_neg_one,
A(i+1,0), &lda,
Y(i,0), &ldy, &c_one,
A(i+1,i), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &i, Y(i,0), &ldy );
#endif
blasf77_sgemv( "No transpose", &m_i1, &i1, &c_neg_one,
X(i+1,0), &ldx,
A(0,i), &ione, &c_one,
A(i+1,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+2:m,i) */
alpha = *A(i+1,i);
lapackf77_slarfg( &m_i1, &alpha, A(min(i+2,m-1),i), &ione, &tauq[i] );
e[i] = MAGMA_S_REAL( alpha );
*A(i+1,i) = c_one;
/* Compute Y(i+1:n,i) */
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_ssetvector( m_i1,
A(i+1,i), 1,
dA(i+1,i), 1, queue );
// 2. Multiply ---------------------------------------------
magma_sgemv( MagmaConjTrans, m_i1, n_i1, c_one,
dA(i+1,i+1), ldda,
dA(i+1,i), ione, c_zero,
dY(i+1,i), ione, queue );
// 3. Get the result back ----------------------------------
magma_sgetmatrix_async( n_i1, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, queue );
blasf77_sgemv( MagmaConjTransStr, &m_i1, &i, &c_one,
A(i+1,0), &lda,
A(i+1,i), &ione, &c_zero,
Y(0,i), &ione );
blasf77_sgemv( "No transpose", &n_i1, &i, &c_neg_one,
Y(i+1,0), &ldy,
Y(0,i), &ione, &c_zero,
work, &ione );
blasf77_sgemv( MagmaConjTransStr, &m_i1, &i1, &c_one,
X(i+1,0), &ldx,
A(i+1,i), &ione, &c_zero,
Y(0,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( queue );
if (i != 0) {
blasf77_saxpy( &n_i1, &c_one, work, &ione, Y(i+1,i), &ione );
}
blasf77_sgemv( MagmaConjTransStr, &i1, &n_i1, &c_neg_one,
A(0,i+1), &lda,
Y(0,i), &ione, &c_one,
Y(i+1,i), &ione );
blasf77_sscal( &n_i1, &tauq[i], Y(i+1,i), &ione );
}
#ifdef COMPLEX
else {
lapackf77_slacgv( &n_i, A(i,i), &lda );
magma_ssetvector( n_i,
A(i,i), lda,
dA(i,i), ldda, queue );
}
#endif
}
}
return info;
} /* magma_slabrd_gpu */
/**
Purpose
-------
SLABRD reduces the first NB rows and columns of a real general
m by n matrix A to upper or lower bidiagonal form by an orthogonal
transformation Q' * A * P, and returns the matrices X and Y which
are needed to apply the transformation to the unreduced part of A.
If m >= n, A is reduced to upper bidiagonal form; if m < n, to lower
bidiagonal form.
This is an auxiliary routine called by SGEBRD.
Arguments
---------
@param[in]
m INTEGER
The number of rows in the matrix A.
@param[in]
n INTEGER
The number of columns in the matrix A.
@param[in]
nb INTEGER
The number of leading rows and columns of A to be reduced.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the m by n general matrix to be reduced.
On exit, the first NB rows and columns of the matrix are
overwritten; the rest of the array is unchanged.
If m >= n, elements on and below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors; and
elements above the diagonal in the first NB rows, with the
array TAUP, represent the orthogonal matrix P as a product
of elementary reflectors.
\n
If m < n, elements below the diagonal in the first NB
columns, with the array TAUQ, represent the orthogonal
matrix Q as a product of elementary reflectors, and
elements on and above the diagonal in the first NB rows,
with the array TAUP, represent the orthogonal matrix P as
a product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[in,out]
dA REAL array, dimension (LDDA,N)
Copy of A on GPU.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
d REAL array, dimension (NB)
The diagonal elements of the first NB rows and columns of
the reduced matrix. D(i) = A(i,i).
@param[out]
e REAL array, dimension (NB)
The off-diagonal elements of the first NB rows and columns of
the reduced matrix.
@param[out]
tauq REAL array dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
@param[out]
taup REAL array, dimension (NB)
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
@param[out]
X REAL array, dimension (LDX,NB)
The m-by-nb matrix X required to update the unreduced part
of A.
@param[in]
ldx INTEGER
The leading dimension of the array X. LDX >= M.
@param[out]
dX REAL array, dimension (LDDX,NB)
Copy of X on GPU.
@param[in]
lddx INTEGER
The leading dimension of the array dX. LDDX >= M.
@param[out]
Y REAL array, dimension (LDY,NB)
The n-by-nb matrix Y required to update the unreduced part
of A.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[out]
dY REAL array, dimension (LDDY,NB)
Copy of Y on GPU.
@param[in]
lddy INTEGER
The leading dimension of the array dY. LDDY >= N.
Further Details
---------------
The matrices Q and P are represented as products of elementary
reflectors:
Q = H(1) H(2) . . . H(nb) and P = G(1) G(2) . . . G(nb)
Each H(i) and G(i) has the form:
H(i) = I - tauq * v * v' and G(i) = I - taup * u * u'
where tauq and taup are real scalars, and v and u are real vectors.
If m >= n, v(1:i-1) = 0, v(i) = 1, and v(i:m) is stored on exit in
A(i:m,i); u(1:i) = 0, u(i+1) = 1, and u(i+1:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
If m < n, v(1:i) = 0, v(i+1) = 1, and v(i+1:m) is stored on exit in
A(i+2:m,i); u(1:i-1) = 0, u(i) = 1, and u(i:n) is stored on exit in
A(i,i+1:n); tauq is stored in TAUQ(i) and taup in TAUP(i).
The elements of the vectors v and u together form the m-by-nb matrix
V and the nb-by-n matrix U' which are needed, with X and Y, to apply
the transformation to the unreduced part of the matrix, using a block
update of the form: A := A - V*Y' - X*U'.
The contents of A on exit are illustrated by the following examples
with nb = 2:
@verbatim
m = 6 and n = 5 (m > n): m = 5 and n = 6 (m < n):
( 1 1 u1 u1 u1 ) ( 1 u1 u1 u1 u1 u1 )
( v1 1 1 u2 u2 ) ( 1 1 u2 u2 u2 u2 )
( v1 v2 a a a ) ( v1 1 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a ) ( v1 v2 a a a a )
( v1 v2 a a a )
@endverbatim
where a denotes an element of the original matrix which is unchanged,
vi denotes an element of the vector defining H(i), and ui an element
of the vector defining G(i).
@ingroup magma_sgesvd_aux
********************************************************************/
extern "C" magma_int_t
magma_slabrd_gpu( magma_int_t m, magma_int_t n, magma_int_t nb,
float *A, magma_int_t lda,
float *dA, magma_int_t ldda,
float *d, float *e, float *tauq, float *taup,
float *X, magma_int_t ldx,
float *dX, magma_int_t lddx,
float *Y, magma_int_t ldy,
float *dY, magma_int_t lddy)
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
#define X(i_,j_) (X + (i_) + (j_)*ldx)
#define Y(i_,j_) (Y + (i_) + (j_)*ldy)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dY(i_,j_) (dY + (i_) + (j_)*lddy)
#define dX(i_,j_) (dX + (i_) + (j_)*lddx)
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float c_zero = MAGMA_S_ZERO;
magma_int_t ione = 1;
magma_int_t i__2, i__3;
magma_int_t i;
float alpha;
A -= 1 + lda;
X -= 1 + ldx;
dX -= 1 + lddx;
Y -= 1 + ldy;
dY -= 1 + lddy;
--d;
--e;
--tauq;
--taup;
/* Quick return if possible */
magma_int_t info = 0;
if (m <= 0 || n <= 0) {
return info;
}
float *f;
magma_queue_t stream;
magma_queue_create( &stream );
magma_smalloc_cpu( &f, max(n,m) );
if ( f == NULL ) {
info = MAGMA_ERR_HOST_ALLOC;
return info;
}
if (m >= n) {
/* Reduce to upper bidiagonal form */
for (i = 1; i <= nb; ++i) {
/* Update A(i:m,i) */
i__2 = m - i + 1;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i,1), &lda,
Y(i,1), &ldy, &c_one,
A(i,i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i,1), &ldx,
A(1,i), &ione, &c_one,
A(i,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+1:m,i) */
alpha = *A(i,i);
i__2 = m - i + 1;
i__3 = i + 1;
lapackf77_slarfg( &i__2, &alpha, A(min(i__3,m),i), &ione, &tauq[i] );
d[i] = MAGMA_S_REAL( alpha );
if (i < n) {
*A(i,i) = c_one;
/* Compute Y(i+1:n,i) */
i__2 = m - i + 1;
i__3 = n - i;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_ssetvector( i__2,
A(i,i), 1,
dA(i-1,i-1), 1 );
// 2. Multiply ---------------------------------------------
magma_sgemv( MagmaConjTrans, i__2, i__3, c_one,
dA(i-1,i), ldda,
dA(i-1,i-1), ione, c_zero,
dY(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_sgetmatrix_async( i__3, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, stream );
i__2 = m - i + 1;
i__3 = i - 1;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
A(i,1), &lda,
A(i,i), &ione, &c_zero,
Y(1,i), &ione );
i__2 = n - i;
i__3 = i - 1;
blasf77_sgemv( "N", &i__2, &i__3, &c_neg_one,
Y(i+1,1), &ldy,
Y(1,i), &ione, &c_zero,
f, &ione );
i__2 = m - i + 1;
i__3 = i - 1;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
X(i,1), &ldx,
A(i,i), &ione, &c_zero,
Y(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__3 != 0) {
i__2 = n - i;
blasf77_saxpy( &i__2, &c_one, f, &ione, Y(i+1,i), &ione );
}
i__2 = i - 1;
i__3 = n - i;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
A(1,i+1), &lda,
Y(1,i), &ione, &c_one,
Y(i+1,i), &ione );
i__2 = n - i;
blasf77_sscal( &i__2, &tauq[i], Y(i+1,i), &ione );
/* Update A(i,i+1:n) */
i__2 = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__2, A(i,i+1), &lda );
lapackf77_slacgv( &i, A(i,1), &lda );
#endif
blasf77_sgemv( "No transpose", &i__2, &i, &c_neg_one,
Y(i+1,1), &ldy,
A(i,1), &lda, &c_one,
A(i,i+1), &lda );
i__2 = i - 1;
i__3 = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i, A(i,1), &lda );
lapackf77_slacgv( &i__2, X(i,1), &ldx );
#endif
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
A(1,i+1), &lda,
X(i,1), &ldx, &c_one,
A(i,i+1), &lda );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__2, X(i,1), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+2:n) */
i__2 = n - i;
i__3 = i + 2;
alpha = *A(i,i+1);
lapackf77_slarfg( &i__2, &alpha, A(i,min(i__3,n)), &lda, &taup[i] );
e[i] = MAGMA_S_REAL( alpha );
*A(i,i+1) = c_one;
/* Compute X(i+1:m,i) */
i__2 = m - i;
i__3 = n - i;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_ssetvector( i__3,
A(i,i+1), lda,
dA(i-1,i), ldda );
// 2. Multiply ---------------------------------------------
//magma_scopy( i__3, dA(i-1,i), ldda, dY(1,1), 1 );
magma_sgemv( MagmaNoTrans, i__2, i__3, c_one,
dA(i,i), ldda,
dA(i-1,i), ldda,
//dY(1,1), 1,
c_zero,
dX(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_sgetmatrix_async( i__2, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, stream );
i__2 = n - i;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i, &c_one,
Y(i+1,1), &ldy,
A(i,i+1), &lda, &c_zero,
X(1,i), &ione );
i__2 = m - i;
blasf77_sgemv( "N", &i__2, &i, &c_neg_one,
A(i+1,1), &lda,
X(1,i), &ione, &c_zero,
f, &ione );
i__2 = i - 1;
i__3 = n - i;
blasf77_sgemv( "N", &i__2, &i__3, &c_one,
A(1,i+1), &lda,
A(i,i+1), &lda, &c_zero,
X(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i != 0) {
i__2 = m - i;
blasf77_saxpy( &i__2, &c_one, f, &ione, X(i+1,i), &ione );
}
i__2 = m - i;
i__3 = i - 1;
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i+1,1), &ldx,
X(1,i), &ione, &c_one,
X(i+1,i), &ione );
i__2 = m - i;
blasf77_sscal( &i__2, &taup[i], X(i+1,i), &ione );
#if defined(PRECISION_z) || defined(PRECISION_c)
i__2 = n - i;
lapackf77_slacgv( &i__2, A(i,i+1), &lda );
// 4. Send the block reflector A(i+1:m,i) to the GPU after SLACGV()
magma_ssetvector( i__2,
A(i,i+1), lda,
dA(i-1,i), ldda );
#endif
}
}
}
else {
/* Reduce to lower bidiagonal form */
for (i = 1; i <= nb; ++i) {
/* Update A(i,i:n) */
i__2 = n - i + 1;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__2, A(i,i), &lda );
lapackf77_slacgv( &i__3, A(i,1), &lda );
#endif
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
Y(i,1), &ldy,
A(i,1), &lda, &c_one,
A(i,i), &lda );
i__2 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__3, A(i,1), &lda );
lapackf77_slacgv( &i__3, X(i,1), &ldx );
#endif
i__3 = n - i + 1;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_neg_one,
A(1,i), &lda,
X(i,1), &ldx, &c_one,
A(i,i), &lda );
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__2, X(i,1), &ldx );
#endif
/* Generate reflection P(i) to annihilate A(i,i+1:n) */
i__2 = n - i + 1;
i__3 = i + 1;
alpha = *A(i,i);
lapackf77_slarfg( &i__2, &alpha, A(i,min(i__3,n)), &lda, &taup[i] );
d[i] = MAGMA_S_REAL( alpha );
if (i < m) {
*A(i,i) = c_one;
/* Compute X(i+1:m,i) */
i__2 = m - i;
i__3 = n - i + 1;
// 1. Send the block reflector A(i,i+1:n) to the GPU ------
magma_ssetvector( i__3,
A(i,i), lda,
dA(i-1,i-1), ldda );
// 2. Multiply ---------------------------------------------
//magma_scopy( i__3, dA(i-1,i-1), ldda, dY(1,1), 1 );
magma_sgemv( MagmaNoTrans, i__2, i__3, c_one,
dA(i,i-1), ldda,
dA(i-1,i-1), ldda,
//dY(1,1), 1,
c_zero,
dX(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_sgetmatrix_async( i__2, 1,
dX(i+1,i), lddx,
X(i+1,i), ldx, stream );
i__2 = n - i + 1;
i__3 = i - 1;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
Y(i,1), &ldy,
A(i,i), &lda, &c_zero,
X(1,i), &ione );
i__2 = m - i;
i__3 = i - 1;
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i+1,1), &lda,
X(1,i), &ione, &c_zero,
f, &ione );
i__2 = i - 1;
i__3 = n - i + 1;
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_one,
A(1,i), &lda,
A(i,i), &lda, &c_zero,
X(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__2 != 0) {
i__3 = m - i;
blasf77_saxpy( &i__3, &c_one, f, &ione, X(i+1,i), &ione );
}
i__2 = m - i;
i__3 = i - 1;
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
X(i+1,1), &ldx,
X(1,i), &ione, &c_one,
X(i+1,i), &ione );
i__2 = m - i;
blasf77_sscal( &i__2, &taup[i], X(i+1,i), &ione );
i__2 = n - i + 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__2, A(i,i), &lda );
magma_ssetvector( i__2,
A(i,i), lda,
dA(i-1,i-1), ldda );
#endif
/* Update A(i+1:m,i) */
i__2 = m - i;
i__3 = i - 1;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
A(i+1,1), &lda,
Y(i,1), &ldy, &c_one,
A(i+1,i), &ione );
i__2 = m - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_slacgv( &i__3, Y(i,1), &ldy );
#endif
blasf77_sgemv( "No transpose", &i__2, &i, &c_neg_one,
X(i+1,1), &ldx,
A(1,i), &ione, &c_one,
A(i+1,i), &ione );
/* Generate reflection Q(i) to annihilate A(i+2:m,i) */
i__2 = m - i;
i__3 = i + 2;
alpha = *A(i+1,i);
lapackf77_slarfg( &i__2, &alpha, A(min(i__3,m),i), &ione, &tauq[i] );
e[i] = MAGMA_S_REAL( alpha );
*A(i+1,i) = c_one;
/* Compute Y(i+1:n,i) */
i__2 = m - i;
i__3 = n - i;
// 1. Send the block reflector A(i+1:m,i) to the GPU ------
magma_ssetvector( i__2,
A(i+1,i), 1,
dA(i,i-1), 1 );
// 2. Multiply ---------------------------------------------
magma_sgemv( MagmaConjTrans, i__2, i__3, c_one,
dA(i,i), ldda,
dA(i,i-1), ione, c_zero,
dY(i+1,i), ione );
// 3. Put the result back ----------------------------------
magma_sgetmatrix_async( i__3, 1,
dY(i+1,i), lddy,
Y(i+1,i), ldy, stream );
i__2 = m - i;
i__3 = i - 1;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i__3, &c_one,
A(i+1,1), &lda,
A(i+1,i), &ione, &c_zero,
Y(1,i), &ione );
i__2 = n - i;
i__3 = i - 1;
blasf77_sgemv( "No transpose", &i__2, &i__3, &c_neg_one,
Y(i+1,1), &ldy,
Y(1,i), &ione, &c_zero,
f, &ione );
i__2 = m - i;
blasf77_sgemv( MagmaConjTransStr, &i__2, &i, &c_one,
X(i+1,1), &ldx,
A(i+1,i), &ione, &c_zero,
Y(1,i), &ione );
// 4. Sync to make sure the result is back ----------------
magma_queue_sync( stream );
if (i__3 != 0) {
i__2 = n - i;
blasf77_saxpy( &i__2, &c_one, f, &ione, Y(i+1,i), &ione );
}
i__2 = n - i;
blasf77_sgemv( MagmaConjTransStr, &i, &i__2, &c_neg_one,
A(1,i+1), &lda,
Y(1,i), &ione, &c_one,
Y(i+1,i), &ione );
i__2 = n - i;
blasf77_sscal( &i__2, &tauq[i], Y(i+1,i), &ione );
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else {
i__2 = n - i + 1;
lapackf77_slacgv( &i__2, A(i,i), &lda );
magma_ssetvector( i__2,
A(i,i), lda,
dA(i-1,i-1), ldda );
}
#endif
}
}
magma_queue_destroy( stream );
magma_free_cpu( f );
return info;
} /* magma_slabrd_gpu */
extern "C" magma_int_t
magma_slaqps(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
float *A, magma_int_t lda,
float *dA, magma_int_t ldda,
magma_int_t *jpvt, float *tau, float *vn1, float *vn2,
float *auxv,
float *F, magma_int_t ldf,
float *dF, magma_int_t lddf)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
SLAQPS computes a step of QR factorization with column pivoting
of a real M-by-N matrix A by using Blas-3. It tries to factorize
NB columns from A starting from the row OFFSET+1, and updates all
of the matrix with Blas-3 xGEMM.
In some cases, due to catastrophic cancellations, it cannot
factorize NB columns. Hence, the actual number of factorized
columns is returned in KB.
Block A(1:OFFSET,1:N) is accordingly pivoted, but not factorized.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0
OFFSET (input) INTEGER
The number of rows of A that have been factorized in
previous steps.
NB (input) INTEGER
The number of columns to factorize.
KB (output) INTEGER
The number of columns actually factorized.
A (input/output) REAL array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, block A(OFFSET+1:M,1:KB) is the triangular
factor obtained and block A(1:OFFSET,1:N) has been
accordingly pivoted, but no factorized.
The rest of the matrix, block A(OFFSET+1:M,KB+1:N) has
been updated.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
JPVT (input/output) INTEGER array, dimension (N)
JPVT(I) = K <==> Column K of the full matrix A has been
permuted into position I in AP.
TAU (output) REAL array, dimension (KB)
The scalar factors of the elementary reflectors.
VN1 (input/output) DOUBLE PRECISION array, dimension (N)
The vector with the partial column norms.
VN2 (input/output) DOUBLE PRECISION array, dimension (N)
The vector with the exact column norms.
AUXV (input/output) REAL array, dimension (NB)
Auxiliar vector.
F (input/output) REAL array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
LDF (input) INTEGER
The leading dimension of the array F. LDF >= max(1,N).
===================================================================== */
#define A(i, j) (A + (i) + (j)*(lda ))
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define F(i, j) (F + (i) + (j)*(ldf ))
#define dF(i, j) (dF + (i) + (j)*(lddf))
float c_zero = MAGMA_S_MAKE( 0.,0.);
float c_one = MAGMA_S_MAKE( 1.,0.);
float c_neg_one = MAGMA_S_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
float d__1;
float z__1;
magma_int_t j, k, rk;
float Akk;
magma_int_t pvt;
float temp, temp2, tol3z;
magma_int_t itemp;
magma_int_t lsticc;
magma_int_t lastrk;
lastrk = min( m, n + offset );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
magma_queue_t stream;
magma_queue_create( &stream );
lsticc = 0;
k = 0;
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// Fortran: pvt, k, isamax are all 1-based; subtract 1 from k.
// C: pvt, k, isamax are all 0-based; don't subtract 1.
pvt = k + cblas_isamax( n-k, &vn1[k], ione );
if (pvt != k) {
if (pvt >= nb) {
/* 1. Start copy from GPU */
magma_sgetmatrix_async( m - offset - nb, 1,
dA(offset + nb, pvt), ldda,
A (offset + nb, pvt), lda, stream );
}
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
blasf77_sswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
vn1[pvt] = vn1[k];
vn2[pvt] = vn2[k];
if (pvt < nb){
/* no need of transfer if pivot is within the panel */
blasf77_sswap( &m, A(0, pvt), &ione, A(0, k), &ione );
}
else {
/* 1. Finish copy from GPU */
magma_queue_sync( stream );
/* 2. Swap as usual on CPU */
blasf77_sswap(&m, A(0, pvt), &ione, A(0, k), &ione);
/* 3. Restore the GPU */
magma_ssetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}
}
/* Apply previous Householder reflectors to column K:
A(RK:M,K) := A(RK:M,K) - A(RK:M,1:K-1)*F(K,1:K-1)'.
Optimization: multiply with beta=0; wait for vector and subtract */
if (k > 0) {
#if defined(PRECISION_c) || defined(PRECISION_z)
for (j = 0; j < k; ++j){
*F(k,j) = MAGMA_S_CNJG( *F(k,j) );
}
#endif
i__1 = m - rk;
i__2 = k;
blasf77_sgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );
#if defined(PRECISION_c) || defined(PRECISION_z)
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_S_CNJG( *F(k,j) );
}
#endif
}
/* Generate elementary reflector H(k). */
if (rk < m-1) {
i__1 = m - rk;
lapackf77_slarfg( &i__1, A(rk, k), A(rk + 1, k), &ione, &tau[k] );
} else {
lapackf77_slarfg( &ione, A(rk, k), A(rk, k), &ione, &tau[k] );
}
Akk = *A(rk, k);
*A(rk, k) = c_one;
/* Compute Kth column of F:
Compute F(K+1:N,K) := tau(K)*A(RK:M,K+1:N)'*A(RK:M,K) on the GPU */
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Send the vector to the GPU */
magma_ssetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL SGEMV( 'Conjugate transpose', M-RK+1, N-K,
// TAU( K ), A( RK, K+1 ), LDA,
// A( RK, K ), 1,
// CZERO, F( K+1, K ), 1 )
magma_int_t i__3 = nb-k-1;
magma_int_t i__4 = i__2 - i__3;
magma_int_t i__5 = nb-k;
magma_sgemv( MagmaTrans, i__1 - i__5, i__2 - i__3,
tau[k], dA(rk +i__5, k+1+i__3), ldda,
dA(rk +i__5, k ), ione,
c_zero, dF(k+1+i__3, k ), ione );
magma_sgetmatrix_async( i__2-i__3, 1,
dF(k + 1 +i__3, k), i__2,
F (k + 1 +i__3, k), i__2, stream );
blasf77_sgemv( MagmaTransStr, &i__1, &i__3,
&tau[k], A(rk, k+1), &lda,
A(rk, k ), &ione,
&c_zero, F(k+1, k ), &ione );
magma_queue_sync( stream );
blasf77_sgemv( MagmaTransStr, &i__5, &i__4,
&tau[k], A(rk, k+1+i__3), &lda,
A(rk, k ), &ione,
&c_one, F(k+1+i__3, k ), &ione );
}
/* Padding F(1:K,K) with zeros. */
for (j = 0; j < k; ++j) {
*F(j, k) = c_zero;
}
/* Incremental updating of F:
F(1:N,K) := F(1:N,K) - tau(K)*F(1:N,1:K-1)*A(RK:M,1:K-1)'*A(RK:M,K). */
if (k > 0) {
i__1 = m - rk;
i__2 = k;
z__1 = MAGMA_S_NEGATE( tau[k] );
blasf77_sgemv( MagmaTransStr, &i__1, &i__2,
&z__1, A(rk, 0), &lda,
A(rk, k), &ione,
&c_zero, auxv, &ione );
i__1 = k;
blasf77_sgemv( MagmaNoTransStr, &n, &i__1,
&c_one, F(0,0), &ldf,
auxv, &ione,
&c_one, F(0,k), &ione );
}
/* Optimization: On the last iteration start sending F back to the GPU */
/* Update the current row of A:
A(RK,K+1:N) := A(RK,K+1:N) - A(RK,1:K)*F(K+1:N,1:K)'. */
if (k < n-1) {
i__1 = n - k - 1;
i__2 = k + 1;
blasf77_sgemm( MagmaNoTransStr, MagmaTransStr, &ione, &i__1, &i__2,
&c_neg_one, A(rk, 0 ), &lda,
F(k+1,0 ), &ldf,
&c_one, A(rk, k+1), &lda );
}
/* Update partial column norms. */
if (rk < lastrk) {
for (j = k + 1; j < n; ++j) {
if (vn1[j] != 0.) {
/* NOTE: The following 4 lines follow from the analysis in
Lapack Working Note 176. */
temp = MAGMA_S_ABS( *A(rk,j) ) / vn1[j];
temp = max( 0., ((1. + temp) * (1. - temp)) );
d__1 = vn1[j] / vn2[j];
temp2 = temp * (d__1 * d__1);
if (temp2 <= tol3z) {
vn2[j] = (float) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_ssqrt(temp);
}
}
}
}
*A(rk, k) = Akk;
++k;
}
// leave k as the last column done
--k;
*kb = k + 1;
rk = offset + *kb - 1;
/* Apply the block reflector to the rest of the matrix:
A(OFFSET+KB+1:M,KB+1:N) := A(OFFSET+KB+1:M,KB+1:N) - A(OFFSET+KB+1:M,1:KB)*F(KB+1:N,1:KB)' */
if (*kb < min(n, m - offset)) {
i__1 = m - rk - 1;
i__2 = n - *kb;
/* Send F to the GPU */
magma_ssetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );
magma_sgemm( MagmaNoTrans, MagmaTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), i__2,
c_one, dA(rk+1, *kb), ldda );
}
/* Recomputation of difficult columns. */
while( lsticc > 0 ) {
itemp = (magma_int_t)(vn2[lsticc] >= 0. ? floor(vn2[lsticc] + .5) : -floor(.5 - vn2[lsticc]));
i__1 = m - rk - 1;
if (lsticc <= nb)
vn1[lsticc] = cblas_snrm2(i__1, A(rk + 1, lsticc), ione);
else {
/* Where is the data, CPU or GPU ? */
float r1, r2;
r1 = cblas_snrm2(nb-k, A(rk + 1, lsticc), ione);
r2 = magma_snrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
//vn1[lsticc] = magma_snrm2(i__1, dA(rk + 1, lsticc), ione);
vn1[lsticc] = magma_ssqrt(r1*r1+r2*r2);
}
/* NOTE: The computation of VN1( LSTICC ) relies on the fact that
SNRM2 does not fail on vectors with norm below the value of SQRT(SLAMCH('S')) */
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;
}
magma_queue_destroy( stream );
return MAGMA_SUCCESS;
} /* magma_slaqps */
/**
Purpose
-------
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by SGEQRF_GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in]
dA REAL array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
SGEQRF_GPU in the first n columns of its array argument A.
@param[in]
ldda INTEGER
The leading dimension of the array A, LDDA >= M.
@param[in]
tau REAL array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_SGEQRF_GPU.
@param[in,out]
dB REAL array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in]
dT REAL array that is the output (the 6th argument)
of magma_sgeqrf_gpu of size
2*MIN(M, N)*NB + ((N+31)/32*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
((N+31)/32*32 )* MAX(NB, NRHS).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) REAL array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK,
LWORK >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by magma_get_sgeqrf_nb( M ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the WORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_sgels_comp
********************************************************************/
extern "C" magma_int_t
magma_sgeqrs_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
float *dA, magma_int_t ldda,
float *tau, float *dT,
float *dB, magma_int_t lddb,
float *hwork, magma_int_t lwork,
magma_int_t *info)
{
#define dA(a_1,a_2) (dA + (a_2)*(ldda) + (a_1))
#define dT(a_1) (dT + (lddwork+(a_1))*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float *dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_sgeqrf_nb(m);
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
int lquery = (lwork == -1);
hwork[0] = MAGMA_S_MAKE( (float)lwkopt, 0. );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -9;
else if (lwork < lwkopt && ! lquery)
*info = -11;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
hwork[0] = c_one;
return *info;
}
/* B := Q' * B */
magma_sormqr_gpu( MagmaLeft, MagmaTrans,
m, nrhs, n,
dA(0,0), ldda, tau,
dB, lddb, hwork, lwork, dT, nb, info );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= k;
if (nb < k)
dwork = dT+2*lddwork*nb;
else
dwork = dT;
// To do: Why did we have this line originally; seems to be a bug (Stan)?
// dwork = dT;
i = (k-1)/nb * nb;
ib = n-i;
rows = m-i;
// TODO: this assumes that, on exit from magma_sormqr_gpu, hwork contains
// the last block of A and B (i.e., C in sormqr). This should be fixed.
// Seems this data should already be on the GPU, so could switch to
// magma_strsm and drop the ssetmatrix.
if ( nrhs == 1 ) {
blasf77_strsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_strsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_ssetmatrix( ib, nrhs, hwork+rows*ib, rows, dwork+i, lddwork );
// update c
if (nrhs == 1)
magma_sgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1);
else
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb);
int start = i-nb;
if (nb < k) {
for (i = start; i >= 0; i -= nb) {
ib = min(k-i, nb);
rows = m -i;
if (i + ib < n) {
if (nrhs == 1) {
magma_sgemv( MagmaNoTrans, ib, ib,
c_one, dT(i), ib,
dB+i, 1,
c_zero, dwork+i, 1);
magma_sgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1);
}
else {
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
ib, nrhs, ib,
c_one, dT(i), ib,
dB+i, lddb,
c_zero, dwork+i, lddwork);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb);
}
}
}
}
magma_scopymatrix( (n), nrhs,
dwork, lddwork,
dB, lddb );
return *info;
}
/**
Purpose
-------
SLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by SGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
A REAL array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T REAL array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y REAL array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
@verbatim
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_sgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_slahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
float *A, magma_int_t lda,
float *tau,
float *T, magma_int_t ldt,
float *Y, magma_int_t ldy,
struct sgehrd_data *data )
{
#define A( i, j ) ( A + (i) + (j)*lda)
#define Y( i, j ) ( Y + (i) + (j)*ldy)
#define T( i, j ) ( T + (i) + (j)*ldt)
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float tmp;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t ione = 1;
magma_int_t d, dki1, dn, nblocks, gblock, lblock, lgid;
magma_int_t n_k_i_1, n_k;
float scale;
magma_int_t i;
float ei = MAGMA_S_ZERO;
magma_int_t info_data = 0;
magma_int_t *info = &info_data;
if (n < 0) {
*info = -1;
} else if (k < 0 || k >= n) {
*info = -2;
} else if (nb < 1 || nb > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldt < nb) {
*info = -8;
} else if (ldy < max(1,n)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// adjust from 1-based indexing
k -= 1;
// Function Body
if (n <= 1)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// zero out current top block of V on all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablas_slaset( MagmaFull, nb, nb, c_zero, c_zero, dV(d,k,0), ldv, data->queues[d] );
}
// set all Y=0
lapackf77_slaset( "Full", &n, &nb, &c_zero, &c_zero, Y, &ldy );
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Finish applying I - V * T * V' on right
tmp = MAGMA_S_NEGATE( tau[i-1] );
blasf77_saxpy( &n_k, &tmp, Y(k,i-1), &ione, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_scopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_strmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_sgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_strmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_sgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_strmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_saxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_slarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// compute yi = A vi = sum_g A{d} vi{d}
nblocks = (n-1) / nb / ngpu + 1;
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
// dV(k+i+1:n-1, i) = VA(k+i:n, i)
magma_ssetvector_async( n_k_i_1,
A(k+i+1,i), 1,
dV(d, k+i+1, i), 1, data->queues[d] );
// copy column of dV -> dVd, using block cyclic distribution.
// This assumes V and Vd have been padded so that
// a 2D matrix copy doesn't access them out-of-bounds
gblock = k / nb;
lblock = gblock / ngpu;
lgid = gblock % ngpu;
if ( d < lgid ) {
lblock += 1;
}
// treat V as (nb*ngpu) x nblock matrix, and Vd as nb x nblock matrix
magmablas_slacpy( MagmaFull, nb, nblocks-lblock,
dV (d, d*nb + lblock*nb*ngpu, i), nb*ngpu,
dVd(d, 0 + lblock*nb, i), nb, data->queues[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
// dY(k:n, i) = dA(k:n, k+i+1:n) * dV(k+i+1:n, i)
// skip if matrix is empty
// each GPU copies to different temporary vector in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_sgemv( MagmaNoTrans, n-k, dn-dki1,
c_one, dA (d, k, dki1), ldda,
dVd(d, dki1, i), 1,
c_zero, dY (d, k, i), 1, data->queues[d] );
// copy vector to host, storing in column nb+d of Y
// as temporary space (Y has >= nb+ngpu columns)
magma_sgetvector_async( n-k,
dY(d, k, i), 1,
Y(k, nb+d), 1, data->queues[d] );
}
}
// while GPU is doing above Ag*v...
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_S_NEGATE( tau[i] );
blasf77_sgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_strmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// apply reflectors to next column, A(i+1), on right only.
// one axpy will be required to finish this, in the next iteration above
if ( i > 0 && i+1 < nb ) {
// Update next column, A(k:n,i+1), applying Q on right.
// One axpy will be required to finish this, in the next iteration
// above, after yi is computed.
// This updates one more row than LAPACK does (row k),
// making block above panel an even multiple of nb.
// Use last column of T as workspace, w.
magma_int_t i1 = i+1;
// If real, conjugate row of V, and undo afterwards
#ifdef COMPLEX
lapackf77_slacgv( &i1, A(k+i1,0), &lda );
#endif
// w = T(0:i, 0:i+1) * VA(k+i+1, 0:i+1)'
// T is now rectangular, so we use gemv instead of trmv as in lapack.
blasf77_sgemv( "No trans", &i, &i1,
&c_one, T(0,0), &ldt,
A(k+i1,0), &lda,
&c_zero, T(0,nb-1), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_sgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i1), &ione );
}
// yi = sum_g yi{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( data->queues[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// yi = yi + yi{d}
blasf77_saxpy( &n_k, &c_one, Y(k,nb+d), &ione, Y(k,i), &ione );
}
}
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
// compute Y = Am V = sum_g Am{d} V{d} --- top part, Y(0:k-1,:)
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
// dY(0:k, :) = dA(0:k, k+i+1:n-1) * dV(k+i+1:n-1, :)
// skip if matrix is empty
// each GPU copies to different temporary block in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_sgemm( MagmaNoTrans, MagmaNoTrans, k, nb, dn-dki1,
c_one, dA (d, 0, dki1), ldda,
dVd(d, dki1, 0), ldvd,
c_zero, dY (d, 0, 0), ldda, data->queues[d] );
// copy result to host, storing in columns [nb + nb*d : nb + nb*(d+1)] of Y
// as temporary space (Y has nb + nb*ngpu columns)
magma_sgetmatrix_async( k, nb,
dY(d, 0, 0), ldda,
Y(0,nb+nb*d), ldy, data->queues[d] );
}
}
// Y = sum_g Y{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( 0 );
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// Y = Y + Am V
for( i = 0; i < nb; ++i ) {
blasf77_saxpy( &k, &c_one, Y(0,nb+nb*d+i), &ione, Y(0,i), &ione );
}
}
}
// copy Y and T matrices to GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_ssetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->queues[d] );
magma_ssetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->queues[d] );
}
magma_setdevice( orig_dev );
return *info;
} /* magma_slahr2 */
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
float *A, *B, *C, *C2, *LU;
float *dA, *dB, *dC1, *dC2;
float alpha = MAGMA_S_MAKE( 0.5, 0.1 );
float beta = MAGMA_S_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float 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_err_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 i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
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 = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_smalloc_pinned( &A, size ); assert( err == 0 );
err = magma_smalloc_pinned( &B, size ); assert( err == 0 );
err = magma_smalloc_pinned( &C, size ); assert( err == 0 );
err = magma_smalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_smalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_smalloc( &dA, size ); assert( err == 0 );
err = magma_smalloc( &dB, size ); assert( err == 0 );
err = magma_smalloc( &dC1, size ); assert( err == 0 );
err = magma_smalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_slarnv( &ione, ISEED, &size, A );
lapackf77_slarnv( &ione, ISEED, &size, B );
lapackf77_slarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test SSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_ssetmatrix( m, n, A, ld, dA, ld );
magma_ssetmatrix( m, n, A, ld, dB, ld );
magma_sswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_sswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasSaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_sgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_slange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "sswap diff %.2g\n", error );
// ----- test ISAMAX
// get argmax of column of A
magma_ssetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_isamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIsamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "isamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test SGEMV
// 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_ssetmatrix( m, n, A, ld, dA, ld );
magma_ssetvector( maxn, B, 1, dB, 1 );
magma_ssetvector( maxn, C, 1, dC1, 1 );
magma_ssetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_sgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSgemv( 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] == 'N' ? m : n);
cublasSaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_SGEMV( m, n ) / 1e9;
printf( "sgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test SSYMV
// 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_ssetmatrix( m, m, A, ld, dA, ld );
magma_ssetvector( m, B, 1, dB, 1 );
magma_ssetvector( m, C, 1, dC1, 1 );
magma_ssetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ssymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsymv( 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
cublasSaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYMV( m ) / 1e9;
printf( "ssymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test STRSV
// 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_slacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_sgetrf( &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_ssetmatrix( m, m, LU, ld, dA, ld );
magma_ssetvector( m, C, 1, dC1, 1 );
magma_ssetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_strsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "strsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test SGEMM
// 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] == 'N');
bool ntb = (trans[ib] == 'N');
magma_ssetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_ssetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_sgemm( 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 );
cublasSgemm( trans[ia], 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
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SGEMM( m, n, k ) / 1e9;
printf( "sgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYMM
// 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_ssetmatrix( m, m, A, ld, dA, ld );
magma_ssetmatrix( m, n, B, ld, dB, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssymm( 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 );
cublasSsymm( side[is], 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
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYMM( side[is], m, n ) / 1e9;
printf( "ssymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYRK
// 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_ssetmatrix( n, k, A, ld, dA, ld );
magma_ssetmatrix( n, n, C, ld, dC1, ld );
magma_ssetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsyrk( uplo[iu], 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
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYRK( k, n ) / 1e9;
printf( "ssyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYR2K
// 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] == 'N');
magma_ssetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_ssetmatrix( n, n, C, ld, dC1, ld );
magma_ssetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssyr2k( 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 );
cublasSsyr2k( uplo[iu], 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
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYR2K( k, n ) / 1e9;
printf( "ssyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test STRMM
// 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] == 'L');
magma_ssetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_strmm( 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 );
cublasStrmm( side[is], uplo[iu], trans[it], 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
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRMM( side[is], m, n ) / 1e9;
printf( "strmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test STRSM
// 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] == 'L');
magma_ssetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_strsm( 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 );
cublasStrsm( side[is], uplo[iu], trans[it], 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
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRSM( side[is], m, n ) / 1e9;
printf( "strsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], 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 );
}
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();
return 0;
}
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
float magma_error, dev_error, work[1];
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t M, N, Xm, Ym, lda, sizeA, sizeX, sizeY;
magma_int_t incx = 1;
magma_int_t incy = 1;
float c_neg_one = MAGMA_S_NEG_ONE;
float alpha = MAGMA_S_MAKE( 1.5, -2.3 );
float beta = MAGMA_S_MAKE( -0.6, 0.8 );
float *A, *X, *Y, *Ydev, *Ymagma;
magmaFloat_ptr dA, dX, dY;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("trans = %s\n", lapack_trans_const(opts.transA) );
#ifdef HAVE_CUBLAS
printf(" M N MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
printf(" M N %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
lda = ((M+31)/32)*32;
gflops = FLOPS_SGEMV( M, N ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
Xm = N;
Ym = M;
} else {
Xm = M;
Ym = N;
}
sizeA = lda*N;
sizeX = incx*Xm;
sizeY = incy*Ym;
TESTING_MALLOC_CPU( A, float, sizeA );
TESTING_MALLOC_CPU( X, float, sizeX );
TESTING_MALLOC_CPU( Y, float, sizeY );
TESTING_MALLOC_CPU( Ydev, float, sizeY );
TESTING_MALLOC_CPU( Ymagma, float, sizeY );
TESTING_MALLOC_DEV( dA, float, sizeA );
TESTING_MALLOC_DEV( dX, float, sizeX );
TESTING_MALLOC_DEV( dY, float, sizeY );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &sizeA, A );
lapackf77_slarnv( &ione, ISEED, &sizeX, X );
lapackf77_slarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_ssetmatrix( M, N, A, lda, dA, 0, lda, opts.queue );
magma_ssetvector( Xm, X, incx, dX, 0, incx, opts.queue );
magma_ssetvector( Ym, Y, incy, dY, 0, incy, opts.queue );
#ifdef HAVE_CUBLAS
dev_time = magma_sync_wtime( 0 );
cublasSgemv( opts.handle, cublas_trans_const(opts.transA),
M, N, &alpha, dA, lda, dX, incx, &beta, dY, incy );
dev_time = magma_sync_wtime( 0 ) - dev_time;
#else
dev_time = magma_sync_wtime( opts.queue );
magma_sgemv( opts.transA, M, N,
alpha, dA, 0, lda,
dX, 0, incx,
beta, dY, 0, incy, opts.queue );
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
#endif
dev_perf = gflops / dev_time;
magma_sgetvector( Ym, dY, 0, incy, Ydev, incy, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_ssetvector( Ym, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
magmablas_sgemv( opts.transA, M, N, alpha, dA, lda, dX, incx, beta, dY, incy );
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_sgetvector( Ym, dY, incy, Ymagma, incy );
#endif
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_sgemv( lapack_trans_const(opts.transA), &M, &N,
&alpha, A, &lda,
X, &incx,
&beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
float Anorm = lapackf77_slange( "F", &M, &N, A, &lda, work );
float Xnorm = lapackf77_slange( "F", &Xm, &ione, X, &Xm, work );
blasf77_saxpy( &Ym, &c_neg_one, Y, &incy, Ydev, &incy );
dev_error = lapackf77_slange( "F", &Ym, &ione, Ydev, &Ym, work ) / (Anorm * Xnorm);
#ifdef HAVE_CUBLAS
blasf77_saxpy( &Ym, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_slange( "F", &Ym, &ione, Ymagma, &Ym, work ) / (Anorm * Xnorm);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ydev );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
extern "C" magma_err_t
magma_sgeqrs_gpu(magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaFloat_ptr dA, size_t dA_offset, magma_int_t ldda,
float *tau, magmaFloat_ptr dT, size_t dT_offset,
magmaFloat_ptr dB, size_t dB_offset, magma_int_t lddb,
float *hwork, magma_int_t lwork,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMagma (version 0.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
Solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by SGEQRF_GPU.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. M >= N >= 0.
NRHS (input) INTEGER
The number of columns of the matrix C. NRHS >= 0.
A (input) REAL array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
SGEQRF_GPU in the first n columns of its array argument A.
LDDA (input) INTEGER
The leading dimension of the array A, LDDA >= M.
TAU (input) REAL array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_SGEQRF_GPU.
DB (input/output) REAL array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
DT (input) REAL array that is the output (the 6th argument)
of magma_sgeqrf_gpu of size
2*MIN(M, N)*NB + ((N+31)/32*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
((N+31)/32*32 )* MAX(NB, NRHS).
LDDB (input) INTEGER
The leading dimension of the array DB. LDDB >= M.
HWORK (workspace/output) REAL array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK, LWORK >= max(1,NRHS).
For optimum performance LWORK >= (M-N+NB)*(NRHS + 2*NB), where
NB is the blocksize given by magma_get_sgeqrf_nb( M ).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the WORK array.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
===================================================================== */
#define a_ref(a_1,a_2) dA, (dA_offset + (a_1) + (a_2)*(ldda))
#define d_ref(a_1) dT, (dT_offset + (lddwork+(a_1))*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magmaFloat_ptr dwork;
magma_int_t i, k, lddwork, rows, ib;
magma_int_t ione = 1;
magma_int_t nb = magma_get_sgeqrf_nb(m);
magma_int_t lwkopt = (m-n+nb)*(nrhs+2*nb);
long int lquery = (lwork == -1);
hwork[0] = MAGMA_S_MAKE( (float)lwkopt, 0. );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -8;
else if (lwork < lwkopt && ! lquery)
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
hwork[0] = c_one;
return *info;
}
/* B := Q' * B */
magma_sormqr_gpu( MagmaLeft, MagmaTrans,
m, nrhs, n,
a_ref(0,0), ldda, tau,
dB, dB_offset, lddb, hwork, lwork, dT, dT_offset, nb, info, queue );
if ( *info != 0 ) {
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= k;
int ldtwork;
size_t dwork_offset = 0;
if (nb < k)
{
dwork = dT;
dwork_offset = dT_offset+2*lddwork*nb;
}
else
{
ldtwork = ( 2*k + ((n+31)/32)*32 )*nb;
magma_smalloc( &dwork, ldtwork );
}
// To do: Why did we have this line originally; seems to be a bug (Stan)?
//dwork = dT;
i = (k-1)/nb * nb;
ib = n-i;
rows = m-i;
if ( nrhs == 1 ) {
blasf77_strsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_strsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_ssetmatrix( ib, nrhs, hwork+rows*ib, 0, rows, dwork, dwork_offset+i, lddwork, queue );
// update c
if (nrhs == 1)
magma_sgemv( MagmaNoTrans, i, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, 1,
c_one, dB, dB_offset, 1, queue );
else
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset + i, lddwork,
c_one, dB, dB_offset, lddb, queue );
int start = i-nb;
if (nb < k) {
for (i = start; i >=0; i -= nb) {
ib = min(k-i, nb);
rows = m -i;
if (i + ib < n) {
if (nrhs == 1) {
magma_sgemv( MagmaNoTrans, ib, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, 1,
c_zero, dwork, dwork_offset+i, 1, queue );
magma_sgemv( MagmaNoTrans, i, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, 1,
c_one, dB, dB_offset, 1, queue );
} else {
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
ib, nrhs, ib,
c_one, d_ref(i), ib,
dB, dB_offset+i, lddb,
c_zero, dwork, dwork_offset+i, lddwork, queue );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
i, nrhs, ib,
c_neg_one, a_ref(0, i), ldda,
dwork, dwork_offset+i, lddwork,
c_one, dB, dB_offset, lddb, queue );
}
}
}
}
magma_scopymatrix( (n), nrhs,
dwork, dwork_offset, lddwork,
dB, dB_offset, lddb, queue );
if (nb >= k)
magma_free(dwork);
magma_queue_sync( queue );
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing zdot
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
const float one = MAGMA_S_MAKE(1.0, 0.0);
const float zero = MAGMA_S_MAKE(0.0, 0.0);
float alpha;
TESTING_INIT();
magma_s_matrix a={Magma_CSR}, b={Magma_CSR}, x={Magma_CSR}, y={Magma_CSR}, skp={Magma_CSR};
printf("%%=======================================================================================================================================================================\n");
printf("\n");
printf(" | runtime | GFLOPS\n");
printf("%% n num_vecs | CUDOT CUGEMV MAGMAGEMV MDOT MDGM MDGM_SHFL | CUDOT CUGEMV MAGMAGEMV MDOT MDGM MDGM_SHFL\n");
printf("%%------------------------------------------------------------------------------------------------------------------------------------------------------------------------\n");
printf("\n");
for( magma_int_t num_vecs=1; num_vecs <= 32; num_vecs += 1 ) {
for( magma_int_t n=500000; n < 500001; n += 10000 ) {
int iters = 10;
float computations = (2.* n * iters * num_vecs);
#define ENABLE_TIMER
#ifdef ENABLE_TIMER
real_Double_t mdot1, mdot2, mdgm1, mdgm2, magmagemv1, magmagemv2, cugemv1, cugemv2, cudot1, cudot2;
real_Double_t mdot_time, mdgm_time, mdgmshf_time, magmagemv_time, cugemv_time, cudot_time;
#endif
CHECK( magma_svinit( &a, Magma_DEV, n, num_vecs, one, queue ));
CHECK( magma_svinit( &b, Magma_DEV, n, 1, one, queue ));
CHECK( magma_svinit( &x, Magma_DEV, n, 8, one, queue ));
CHECK( magma_svinit( &y, Magma_DEV, n, 8, one, queue ));
CHECK( magma_svinit( &skp, Magma_DEV, 1, num_vecs, zero, queue ));
// warm up
CHECK( magma_sgemvmdot( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
// CUDOT
#ifdef ENABLE_TIMER
cudot1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
for( int l=0; l<num_vecs; l++){
alpha = magma_sdot( n, a.dval+l*a.num_rows, 1, b.dval, 1, queue );
//cudaDeviceSynchronize();
}
//cudaDeviceSynchronize();
}
#ifdef ENABLE_TIMER
cudot2 = magma_sync_wtime( queue );
cudot_time=cudot2-cudot1;
#endif
// CUGeMV
#ifdef ENABLE_TIMER
cugemv1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
magma_sgemv( MagmaTrans, n, num_vecs, one, a.dval, n, b.dval, 1, zero, skp.dval, 1, queue );
}
#ifdef ENABLE_TIMER
cugemv2 = magma_sync_wtime( queue );
cugemv_time=cugemv2-cugemv1;
#endif
// MAGMAGeMV
#ifdef ENABLE_TIMER
magmagemv1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
magmablas_sgemv( MagmaTrans, n, num_vecs, one, a.dval, n, b.dval, 1, zero, skp.dval, 1, queue );
}
#ifdef ENABLE_TIMER
magmagemv2 = magma_sync_wtime( queue );
magmagemv_time=magmagemv2-magmagemv1;
#endif
// MDOT
#ifdef ENABLE_TIMER
mdot1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
for( int c = 0; c<num_vecs/2; c++ ){
CHECK( magma_smdotc( n, 2, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
for( int c = 0; c<num_vecs%2; c++ ){
CHECK( magma_smdotc( n, 1, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
//h++;
}
#ifdef ENABLE_TIMER
mdot2 = magma_sync_wtime( queue );
mdot_time=mdot2-mdot1;
#endif
// MDGM
#ifdef ENABLE_TIMER
mdgm1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
CHECK( magma_sgemvmdot( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
//h++;
}
#ifdef ENABLE_TIMER
mdgm2 = magma_sync_wtime( queue );
mdgm_time=mdgm2-mdgm1;
#endif
// MDGM_shfl
#ifdef ENABLE_TIMER
mdgm1 = magma_sync_wtime( queue );
#endif
for( int h=0; h < iters; h++) {
CHECK( magma_sgemvmdot_shfl( n, num_vecs, a.dval, b.dval, x.dval, y.dval, skp.dval, queue ));
}
#ifdef ENABLE_TIMER
mdgm2 = magma_sync_wtime( queue );
mdgmshf_time=mdgm2-mdgm1;
#endif
//magma_sprint_gpu(num_vecs,1,skp.dval,num_vecs);
//Chronometry
#ifdef ENABLE_TIMER
printf("%d %d %e %e %e %e %e %e || %e %e %e %e %e %e\n",
int(n), int(num_vecs),
cudot_time/iters,
(cugemv_time)/iters,
(magmagemv_time)/iters,
(mdot_time)/iters,
(mdgm_time)/iters,
(mdgmshf_time)/iters,
computations/(cudot_time*1e9),
computations/(cugemv_time*1e9),
computations/(magmagemv_time*1e9),
computations/(mdot_time*1e9),
computations/(mdgm_time*1e9),
computations/(mdgmshf_time*1e9) );
#endif
magma_smfree(&a, queue );
magma_smfree(&b, queue );
magma_smfree(&x, queue );
magma_smfree(&y, queue );
magma_smfree(&skp, queue );
}
//printf("%%================================================================================================================================================\n");
//printf("\n");
//printf("\n");
}
// use alpha to silence compiler warnings
if ( isnan( real( alpha ))) {
info = -1;
}
cleanup:
magma_queue_destroy( queue );
TESTING_FINALIZE();
return info;
}
