extern "C" magma_int_t
magma_zdtohpo(magma_int_t num_gpus, char *uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb, magma_int_t NB,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *dwork[], magma_int_t ldda,
magma_queue_t stream[][3], magma_int_t *info)
{
magma_int_t k;
if( lapackf77_lsame(uplo, "U") ) {
magma_int_t j, jj, jb, mj;
/* go through each column */
for (j=off_j+NB; j<n; j+=nb) {
jj = (j-off_j)/(nb*num_gpus);
k = ((j-off_j)/nb)%num_gpus;
jb = min(nb, (n-j));
if(j+jb < off_j+m)
mj = (j-off_i)+jb;
else
mj = m;
magma_setdevice(k);
magma_zgetmatrix_async( mj, jb,
dA(k, 0, jj*nb), ldda,
A(off_i, j), lda,
stream[k][0] );
magma_queue_sync( stream[k][0] );
}
} else {
magma_int_t i, ii, ib, ni;
/* go through each row */
for(i=off_i+NB; i<m; i+=nb) {
ii = (i-off_i)/(nb*num_gpus);
k = ((i-off_i)/nb)%num_gpus;
ib = min(nb, (m-i));
if(i+ib < off_i+n)
ni = (i-off_i)+ib;
else
ni = n;
magma_setdevice(k);
magma_zgetmatrix_async( ib, ni,
dA(k, ii*nb, 0), ldda,
A(i, off_j), lda,
stream[k][0] );
magma_queue_sync( stream[k][0] );
}
}
/*for( k=0; k<num_gpus; k++ ) {
magma_setdevice(k);
magma_queue_sync( stream[k][0] );
}*/
magma_setdevice(0);
return *info;
}
// ----------------------------------------------------------------------
// TODO info is unused
extern "C" magma_int_t
magma_zhtodhe(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex_ptr dA[], magma_int_t ldda,
magma_queue_t queues[][10],
magma_int_t *info)
{
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_int_t k;
if (uplo == MagmaLower) {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*ngpu);
k = (j/nb)%ngpu;
jb = min(nb, (n-j));
mj = n-j;
magma_setdevice( k );
magma_zsetmatrix_async( mj, jb,
A(j,j), lda,
dA(k, j, jj*nb), ldda,
queues[k][0] );
}
}
else {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*ngpu);
k = (j/nb)%ngpu;
jb = min(nb, (n-j));
mj = j+jb;
magma_setdevice( k );
magma_zsetmatrix_async( mj, jb,
A(0, j), lda,
dA(k, 0, jj*nb), ldda,
queues[k][0] );
}
}
for( k=0; k < ngpu; k++ ) {
magma_setdevice( k );
magma_queue_sync( queues[k][0] );
}
magma_setdevice( orig_dev );
return *info;
}
extern "C" magma_int_t
magma_shtodhe(magma_int_t num_gpus, magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
float *A, magma_int_t lda,
float **dA, magma_int_t ldda,
magma_queue_t stream[][10], magma_int_t *info)
{
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_int_t k;
if (uplo == MagmaLower) {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*num_gpus);
k = (j/nb)%num_gpus;
jb = min(nb, (n-j));
mj = n-j;
magma_setdevice(k);
magma_ssetmatrix_async( mj, jb,
A(j,j), lda,
dA(k, j, jj*nb), ldda,
stream[k][0] );
}
}
else {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*num_gpus);
k = (j/nb)%num_gpus;
jb = min(nb, (n-j));
mj = j+jb;
magma_setdevice(k);
magma_ssetmatrix_async( mj, jb,
A(0, j), lda,
dA(k, 0, jj*nb), ldda,
stream[k][0] );
}
}
for( k=0; k < num_gpus; k++ ) {
magma_setdevice(k);
magma_queue_sync(stream[k][0]);
}
magma_setdevice( orig_dev );
return *info;
}
Operator &ParNonlinearForm::GetGradient(const Vector &x) const
{
ParFiniteElementSpace *pfes = ParFESpace();
pGrad.Clear();
NonlinearForm::GetGradient(x); // (re)assemble Grad, no b.c.
OperatorHandle dA(pGrad.Type()), Ph(pGrad.Type());
if (fnfi.Size() == 0)
{
dA.MakeSquareBlockDiag(pfes->GetComm(), pfes->GlobalVSize(),
pfes->GetDofOffsets(), Grad);
}
else
{
MFEM_ABORT("TODO: assemble contributions from shared face terms");
}
// TODO - construct Dof_TrueDof_Matrix directly in the pGrad format
Ph.ConvertFrom(pfes->Dof_TrueDof_Matrix());
pGrad.MakePtAP(dA, Ph);
// Impose b.c. on pGrad
OperatorHandle pGrad_e;
pGrad_e.EliminateRowsCols(pGrad, ess_tdof_list);
return *pGrad.Ptr();
}
void cmaxent_fortran( double* xqmc, double* xtau, int32_t len, double xmom1, double ( *xker ) ( const double&, double&, double& ),
double ( *backtrans ) ( double&, double&, double& ), double beta, double* alpha_tot, int32_t n_alpha, int32_t ngamma, double omega_start, double omega_end,
int32_t omega_points, int32_t nsweeps, int32_t nbins, int32_t nwarmup,/* double* u,*/ double* sigma)
{
std::string fr("Aom"); std::string dA("dump_Aom"); std::string ml("max_stoch_log"); std::string energies("energies"); std::string bf("best_fit"); std::string dump("dump");
cmaxent(xqmc, xtau, len, xmom1, xker, backtrans, beta, alpha_tot, n_alpha, ngamma, omega_start, omega_end, omega_points, nsweeps, nbins, nwarmup,
fr, dA, ml, energies, bf, dump, /*u*/NULL, sigma);
}
void testDeviceVector()
{
const int aSize = 64;
std::vector<int> hA(aSize), hB(aSize);
bolt::cl::device_vector<int> dA(aSize), dB(aSize);
for(int i=0; i<aSize; i++) {
hA[i] = hB[i] = dB[i] = dA[i] = i;
};
int hSum = std::inner_product(hA.begin(), hA.end(), hB.begin(), 1);
int sum = bolt::cl::inner_product( dA.begin(), dA.end(),
dB.begin(), 1, bolt::cl::plus<int>(), bolt::cl::multiplies<int>() );
};
void testDeviceVector()
{
const int aSize = 1000;
std::vector<int> hA(aSize);
bolt::cl::device_vector<int> dA(aSize);
for(int i=0; i<aSize; i++) {
hA[i] = i;
dA[i] = i;
};
std::vector<int>::iterator smaxdex = std::max_element(hA.begin(), hA.end());
bolt::cl::device_vector<int>::iterator bmaxdex = bolt::cl::max_element(dA.begin(), dA.end(),bolt::cl::greater<int>());
};
void GLTorus::draw()
{
GLfloat R(m_majorRadius);
GLfloat r(m_minorRadius);
GLfloat dA(m_angle/m_majorSegments);
GLfloat da(TwoPi/m_minorSegments);
GLfloat cosTheta[2], sinTheta[2], cosPhi, sinPhi;
Vec v, n;
for (int i = 0; i < m_majorSegments; ++i) {
cosTheta[0] = cos( i *dA);
cosTheta[1] = cos((i+1)*dA);
sinTheta[0] = sin( i *dA);
sinTheta[1] = sin((i+1)*dA);
glBegin(GL_QUAD_STRIP);
for (int j = 0; j <= m_minorSegments; ++j) {
cosPhi = cos(j*da);
sinPhi = sin(j*da);
for (int k = 0; k <= 1; ++k) {
v.x = (R+r*cosPhi) * cosTheta[k];
v.y = (R+r*cosPhi) * sinTheta[k];
v.z = r*sinPhi;
n.x = R*cosTheta[k];
n.y = R*sinTheta[k];
n.z = 0;
n = (v-n).unit();
glNormal3f(n.x, n.y, n.z);
glVertex3f(v.x, v.y, v.z);
}
}
glEnd();
}
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
double c_neg_one = MAGMA_D_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' };
double *A, *B, *C, *C2, *LU;
double *dA, *dB, *dC1, *dC2;
double alpha = MAGMA_D_MAKE( 0.5, 0.1 );
double beta = MAGMA_D_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_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_dmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_dmalloc( &dA, size ); assert( err == 0 );
err = magma_dmalloc( &dB, size ); assert( err == 0 );
err = magma_dmalloc( &dC1, size ); assert( err == 0 );
err = magma_dmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_dlarnv( &ione, ISEED, &size, B );
lapackf77_dlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test DSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetmatrix( m, n, A, ld, dB, ld );
magma_dswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_dswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasDaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_dgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "dswap diff %.2g\n", error );
// ----- test IDAMAX
// get argmax of column of A
magma_dsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_idamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIdamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "idamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test DGEMV
// 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_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetvector( maxn, B, 1, dB, 1 );
magma_dsetvector( maxn, C, 1, dC1, 1 );
magma_dsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemv( 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);
cublasDaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMV( m, n ) / 1e9;
printf( "dgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DSYMV
// 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_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetvector( m, B, 1, dB, 1 );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymv( 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
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMV( m ) / 1e9;
printf( "dsymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DTRSV
// 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_dlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_dgetrf( &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_dsetmatrix( m, m, LU, ld, dA, ld );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsv( 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
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "dtrsv( %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 DGEMM
// 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_dsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_dsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dgemm( 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 );
cublasDgemm( 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
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMM( m, n, k ) / 1e9;
printf( "dgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYMM
// 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_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetmatrix( m, n, B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsymm( 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 );
cublasDsymm( 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
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMM( side[is], m, n ) / 1e9;
printf( "dsymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYRK
// 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_dsetmatrix( n, k, A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyrk( 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
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYRK( k, n ) / 1e9;
printf( "dsyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYR2K
// 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_dsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyr2k( 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 );
cublasDsyr2k( 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
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYR2K( k, n ) / 1e9;
printf( "dsyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DTRMM
// 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_dsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrmm( 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 );
cublasDtrmm( 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
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRMM( side[is], m, n ) / 1e9;
printf( "dtrmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test DTRSM
// 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_dsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrsm( 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 );
cublasDtrsm( 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
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( side[is], m, n ) / 1e9;
printf( "dtrsm( %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;
}
/**
Purpose
-------
CGEQRF_OOC computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
This is an out-of-core (ooc) version that is similar to magma_cgeqrf but
the difference is that this version can use a GPU even if the matrix
does not fit into the GPU memory at once.
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,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb( M, N ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf_ooc(
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
#define A(i_,j_) ( A + (i_) + (j_)*lda )
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
/* Constants */
const magmaFloatComplex c_one = MAGMA_C_ONE;
/* Local variables */
magmaFloatComplex_ptr dA, dwork;
magma_int_t i, ib, IB, j, min_mn, lddwork, ldda, rows;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
magma_int_t lwkopt = n * nb;
work[0] = magma_cmake_lwork( lwkopt );
bool lquery = (lwork == -1);
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,n) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Check how much memory do we have */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
magma_int_t NB = magma_int_t(0.8*freeMem/m);
NB = (NB / nb) * nb;
if (NB >= n)
return magma_cgeqrf(m, n, A, lda, tau, work, lwork, info);
min_mn = min(m,n);
if (min_mn == 0) {
work[0] = c_one;
return *info;
}
lddwork = magma_roundup( NB, 32 ) + nb;
ldda = magma_roundup( m, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (NB + nb)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magmaFloatComplex_ptr ptr = dA + ldda*NB;
dwork = dA + ldda*(NB + nb);
/* start the main loop over the blocks that fit in the GPU memory */
for (i=0; i < n; i += NB) {
IB = min( n-i, NB );
//printf("Processing %5d columns -- %5d to %5d ... \n", IB, i, i+IB);
/* 1. Copy the next part of the matrix to the GPU */
magma_csetmatrix_async( m, IB,
A(0,i), lda,
dA(0,0), ldda, queues[0] );
magma_queue_sync( queues[0] );
/* 2. Update it with the previous transformations */
for (j=0; j < min(i,min_mn); j += nb) {
ib = min( min_mn-j, nb );
/* Get a panel in ptr. */
// 1. Form the triangular factor of the block reflector
// 2. Send it to the GPU.
// 3. Put 0s in the upper triangular part of V.
// 4. Send V to the GPU in ptr.
// 5. Update the matrix.
// 6. Restore the upper part of V.
rows = m-j;
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(j,j), &lda, tau+j, work, &ib);
magma_csetmatrix_async( ib, ib,
work, ib,
dwork, lddwork, queues[1] );
magma_cpanel_to_q( MagmaUpper, ib, A(j,j), lda, work+ib*ib );
magma_csetmatrix_async( rows, ib,
A(j,j), lda,
ptr, rows, queues[1] );
magma_queue_sync( queues[1] );
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, IB, ib,
ptr, rows, dwork, lddwork,
dA(j, 0), ldda, dwork+ib, lddwork, queues[1] );
magma_cq_to_panel( MagmaUpper, ib, A(j,j), lda, work+ib*ib );
}
/* 3. Do a QR on the current part */
if (i < min_mn)
magma_cgeqrf2_gpu( m-i, IB, dA(i,0), ldda, tau+i, info );
/* 4. Copy the current part back to the CPU */
magma_cgetmatrix_async( m, IB,
dA(0,0), ldda,
A(0,i), lda, queues[0] );
}
magma_queue_sync( queues[0] );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_cgeqrf_ooc */
/**
Purpose
-------
CLATRD2 reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, CLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, CLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by CHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_cheev_aux
********************************************************************/
extern "C" magma_int_t
magma_clatrd2(magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
float *e, magmaFloatComplex *tau,
magmaFloatComplex *W, magma_int_t ldw,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dW, magma_int_t lddw,
magmaFloatComplex *dwork, magma_int_t ldwork)
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
#define dW(i, j) (dW + (j)*lddw + (i))
magma_int_t i;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex value = MAGMA_C_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaFloatComplex alpha;
magmaFloatComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_cmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside clatrd
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_cgemv("No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_clacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_cgemv("No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, A(i, i+1), &ldw);
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_clarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_C_REAL( alpha );
*A(i-1,i) = MAGMA_C_MAKE( 1, 0 );
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_csetvector( i, A(0, i), 1, dA(0, i), 1 );
//#if (GPUSHMEM < 200)
//magma_chemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
// dA(0, i), ione, c_zero, dW(0, iw), ione);
//#else
magmablas_chemv_work(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_cgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_cgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( stream );
if (i < n-1) {
blasf77_cgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_cgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_cgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_cscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_cdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_caxpy(&i, &alpha, A(0, i), &ione,
W(0, iw), &ione);
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, W(i, 0), &ldw);
#endif
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, W(i, 0), &ldw);
lapackf77_clacgv(&i, A(i, 0), &lda);
#endif
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, A(i, 0), &lda);
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_clarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_C_REAL( alpha );
*A(i+1,i) = MAGMA_C_MAKE( 1, 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_csetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
//#if (GPUSHMEM < 200)
//magma_chemv(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
// dW(i+1, i), ione);
//#else
magmablas_chemv_work(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_cgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_cgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_cgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
// 3. Here is where we need it
magma_queue_sync( stream );
if (i != 0)
blasf77_caxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_cscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_cdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_caxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* magma_clatrd */
void magmablas_ssymm_mgpu_com(
magma_side_t side, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
float alpha,
float *dA[], magma_int_t ldda, magma_int_t offset,
float *dB[], magma_int_t lddb,
float beta, float *dC[], magma_int_t lddc,
float *dwork[], magma_int_t dworksiz,
float *C, magma_int_t ldc,
float *work[], magma_int_t worksiz,
magma_int_t ngpu, magma_int_t nb,
magma_queue_t streams[][20], magma_int_t nstream,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10], magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
//printf("####################################################\n");
//printf(" start ssymm \n");
//printf("####################################################\n");
if ( side != MagmaLeft || uplo != MagmaLower ) {
fprintf( stderr, "%s: only Left Lower implemented\n", __func__ );
}
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nstream >= ngpu );
assert( nbevents >= ngpu*ngpu );
float c_one = MAGMA_S_ONE;
float *dwork1[MagmaMaxGPUs];
float *dwork2[MagmaMaxGPUs];
magma_int_t maxgsize = n*m;
magma_int_t lddwork = lddc;
magma_int_t ldwork = m;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev]; // size of dwork1 is n*lddwork
dwork2[dev] = dwork[dev]+n*lddwork; // size of dwork2 is maxgsize*ngpu
}
assert( dworksiz >= (n*lddwork+maxgsize*ngpu) );
assert( worksiz >= (n*ldwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev, devperm, myblk, mycolsize, myblkoffst;
magma_int_t gmaster;
magma_int_t masterdev, lcdev, lccolsize, myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m, nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
cudaMemset(dwork(dev,0,0), 0, (lddwork)*(n)*sizeof(float) );
// put all dC on all dev to 0 except the one which
// hold i==0 because this one has to multiply by beta.
if(dev!=stdev){
cudaMemset(dC(dev,0,0), 0, (lddc)*(n)*sizeof(float) );
}
}
magma_int_t newoffset = offset;
// 1. symmetrize
if(blockoffset>0){
newoffset = offset+fstblksiz; // newoffset is adjusted over nb
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > stdev?1:0);
//printf("STDEV %d voici offset %d remm %d myblockoffset %d siz %d \n", stdev, offset, remm, myblkoffst, fstblksiz);
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
magmablas_ssymmetrize_tiles( MagmaLower, fstblksiz, dA(stdev, offset, myblkoffst*nb+blockoffset), ldda, 1, ngpu*nb, nb );
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t newstdev = (newoffset/nb)%ngpu;
magma_int_t nbblk = remm/nb; // number of block of size nb. if m%nb>0 then a last block exist and is of size ib=m%nb
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-newstdev+ngpu)%ngpu) ? 1:0 );
magma_int_t devperm = (dev-newstdev+ngpu)%ngpu;
magma_int_t nbblkoffst = newoffset/nb;
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
//printf("dev %d devperm %d newoffset %d rowoff %d coloff %d myblk %d \n", dev, devperm, newoffset, newoffset+devperm*nb, myblkoffst*nb, myblk);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magmablas_ssymmetrize_tiles( MagmaLower, nb, dA(dev, newoffset+devperm*nb, myblkoffst*nb), ldda, myblk, ngpu*nb, nb );
if(remm%nb>0){
magma_int_t nblstblks = (nbblk+1)%ngpu;
magma_int_t devlstblk = (nblstblks-1+ngpu)%ngpu;
//printf("==> siz %d devperm %d, devlstblk %d, newoffset+nbblk*nb %d, myblkoffst*nb+ myblk*nb %d\n", remm % nb, devperm, devlstblk, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb);
if(devperm==devlstblk)
magmablas_ssymmetrize( MagmaLower, remm % nb, dA(dev, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb), ldda ); // last partial tile
}
}
/*
magma_int_t siz = m+offset;
float *R;
magma_smalloc_cpu( &R, siz*siz );
// collecte back A
magmablas_sgetmatrix_1D_bcyclic( siz, siz, dA, ldda, R, siz, ngpu, nb );
magma_setdevice( 0 );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
//magma_sgetmatrix( siz, siz, dA[0], ldda, R, siz );
FILE *trace_file;
trace_file = fopen("AJETE/Aafter", "w");
for (int j = 0; j < siz ; j++)
for (int i = 0; i < siz ; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, R[j*siz+i]);
fclose(trace_file);
return;
*/
// ROW GEMM transpose a row and make a gemm with a block
// if only 1 GPU used the ROW GEMM is integrated with the
// COL GEMM (better accuracy observed) and better perf
if(ngpu>1){
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
//magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nbblk = magma_ceildiv(i, nb);
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-stdev+ngpu)%ngpu) ? 1:0 );
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
magma_int_t myrowsize = myblk * nb;
magma_int_t coloffset = myblkoffst*nb;
if(dev==stdev) {
myrowsize = myrowsize -blockoffset;
coloffset = myblkoffst*nb+blockoffset;
}
//printf("ROW GEMM: voici i %d ib %d ioff %d nbblkoffst %d stdev %d dev %d myblk %d myblkoffset %d coloffset %d rowsize %d\n", i, ib, ioff, nbblkoffst, stdev, dev, myblk, myblkoffst, coloffset, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, myrowsize, n, ib,
alpha, dA(dev,ioff,coloffset), ldda,
dB(dev,i,0), lddb,
c_one, dwork(dev,0,0), lddwork );
}
}
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_event_record(redevents[dev][1], streams[dev][1]);
}
}
// COL GEMM
// blockoffset is offset within first block; for subsequent blocks it is 0
if(blockoffset>0){
magma_int_t ib = min( nb-blockoffset, m ); // block size
magma_int_t iblock = (offset / nb) / ngpu; // local block id
magma_int_t di = iblock*nb+blockoffset; // local index in parent matrix
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
//printf("DEV %d COL GEMM first ioff %d di %d m %d n %d ib %d \n", stdev, offset, di, m, n, ib);
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m, n, ib,
alpha, dA(stdev,offset,di), ldda,
dB(stdev,0,0), lddb,
beta, dC(stdev,0,0), lddc );
}
// COL GEMM
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
magma_int_t iblock = (ioff / nb) / ngpu; // local block id
magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t di = iblock*nb; // local index in parent matrix
//printf("DEV %d COL GEMM i %d ioff %d di %d m-i %d n %d ib %d \n", dev, i, ioff, di, m-i, n, ib);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
if(i==0){
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
beta, dC(dev,i,0), lddc );
}else{
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,i,0), lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
// if only 1 GPU is used, do the ROW GEMM
if(ngpu==1){
// NOTE THAT because the COL gemm write dC below the diagonal (i)
// and the ROW GEMM write dC from 0 to diag-1, so they could
// run in parallel on different streams.
//
// NO NO NO because
// it might happen that col finished i and strated i+1 while row still at i
// magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, i, n, ib,
alpha, dA(dev,ioff,offset), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,0,0), lddc );
}
}
if(ngpu>1){
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t nbblkrow = nbblk-1;
magma_int_t devperm = (dev-stdev+ngpu)%ngpu;
magma_int_t myblk = (nbblkrow/ngpu) + (nbblkrow%ngpu > devperm ? 1:0 );
magma_int_t myrowsize = myblk * nb;
if(dev==stdev) {
myrowsize = myrowsize - blockoffset;
}
//printf("blockoffset %d nbblkrow %d devperm %d DEV %d RECEIVING myblk %d myrowsize %d\n", blockoffset, nbblkrow, devperm, dev, myblk, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][1]);
//magma_queue_sync( streams[ dev ][ 1 ] );
// for each dev add the computed ROW block each on its placment with dC
for( magma_int_t blki = 0; blki < myblk; ++blki){
magma_int_t gbblki = (blki*ngpu + devperm)*nb - blockoffset;
magma_int_t lcblki = blki*nb;
magma_int_t ib = nb;// min(nb, m-gbblki);
if(dev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
magmablas_sgeadd(ib, n, c_one,
&dwork[dev][lcblki], lddwork,
&dC[dev][gbblki] , lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
}
}
}
// ===========================================================
// COMMUNICATION ALL_REDUCE_SUM
// ===========================================================
if(ngpu==1){
return;
}
// INITIALIZE COMM
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize, m);
if(mycolsize>0){
gpuisactive[dev] = mycolsize;
if(masterdev==-1) {
masterdev = dev;
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//*******************************
// each GPU send its result
// to its master. The master make
// the addition and then send to
// to the masters of other real
// and receive from the masters of
// other real make the addition
// and broadcast locally the final
// result.
//*******************************
//printf("=======================================================================\n");
//printf(" sending to my master \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU. if I am not the master, then send my result to my master.
// store result on dwork[masterdev][dev*maxgsize]
if(dev!=masterdev){
magma_setdevice( dev );
//printf(" GPU %d sending to my master %d\n", dev, masterdev);
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][0]);
// sending to the master of my real
magma_scopymatrix_async(
m, n,
&dC[dev][0], lddc,
&dwork2[masterdev][maxgsize*dev], m, streams[dev][0] );
magma_event_record(redevents[dev][masterdev], streams[dev][0]);
} // end I am not the masterdev
}// end if mycolsize>0
}// for idev
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//printf("=======================================================================\n");
//printf(" each master do addition of local result and broadcast to other masters \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][0]);
// ========================================
// local addition
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" master %d receiving from %d and adding \n", masterdev, lcdev);
// this is an active GPU of my real.
// wait I received what he send it to me and then do addition.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[lcdev][masterdev]);
magmablas_sgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*lcdev], m,
&dC[masterdev][0] , lddc );
}
}// for l=1:myngpu
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
//
// ========================================
// send to other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//Master has to wait until finish the local addition then send using gmaster stream.
//use stream 0 to make it sequential or stream gmaster to make it parallel.
//Now both re the same.
//printf(" master %d from cmplx %d sending to other master %d on cmplx %d \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ gmaster ], redevents[masterdev][masterdev]);
magma_scopymatrix_async(
m, n,
&dC[masterdev][0], lddc,
&dwork2[gmaster][maxgsize*masterdev], m, streams[masterdev][gmaster] );
magma_event_record(redevents[masterdev][gmaster], streams[masterdev][gmaster]);
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][gmaster]);
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// ========================================
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//printf("=======================================================================\n");
//printf(" each master wait receiving other masters results, do the addition and broadcast locally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// master has to wait until finishing all the send to other masters.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
// ========================================
// addition of results from other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//Master has to wait until receiving from gmaster, then do addition using stream 0
//printf(" master %d from cmplx %d receiving from other master %d on cmplx %d and adding \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[gmaster][masterdev]);
magmablas_sgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*gmaster], m,
&dC[masterdev][0] , lddc );
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
// ========================================
// local broadcast of final results
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
// this is an active GPU of my real.
// wait the previous addition is done maening stream 0 is finished and broadcast sequentially for now.
// to make it parallel put stream lcdev instead of stream 0
//printf(" master %d broadcasting local to %d \n", masterdev, lcdev);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
magma_scopymatrix_async(
m, n,
&dC[masterdev][0], lddc,
&dC[lcdev][0], lddc, streams[masterdev][0] );
magma_event_record(redevents[masterdev][lcdev], streams[masterdev][0]);
}
}// for l=1:myngpu
// ========================================
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if(lccolsize>0){
magma_setdevice( lcdev );
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[lcdev][0]);
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[masterdev][lcdev]);
}
}// for l=1:myngpu
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
//printf("****************************************************\n");
//printf(" finish ssymm \n");
//printf("****************************************************\n");
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}
/**
Purpose
-------
SORGQR generates an M-by-N REAL matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by SGEQRF_GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
dA REAL array A on the GPU device,
dimension (LDDA,N). On entry, the i-th column must contain
the vector which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by SGEQRF_GPU in the first k
columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
ldda INTEGER
The first dimension of the array A. LDDA >= max(1,M).
@param[in]
tau REAL array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by SGEQRF_GPU.
@param[in]
dT REAL work space array on the GPU device,
dimension (MIN(M, N) )*NB.
This must be the 6th argument of magma_sgeqrf_gpu
[ note that if N here is bigger than N in magma_sgeqrf_gpu,
the workspace requirement DT in magma_sgeqrf_gpu must be
as specified in this routine ].
@param[in]
nb INTEGER
This is the block size used in SGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument has an illegal value
@ingroup magma_ssyev_2stage
********************************************************************/
extern "C" magma_int_t
magma_sorgqr_2stage_gpu(magma_int_t m, magma_int_t n, magma_int_t k,
float *dA, magma_int_t ldda,
float *tau, float *dT,
magma_int_t nb, magma_int_t *info)
{
#define dA(a_1,a_2) (dA + (a_2)*(ldda) + (a_1))
#define dT(a_1) (dT + (a_1)*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
magma_int_t i__1, i__2, i__3;
//magma_int_t lwork;
magma_int_t i, ib, ki, kk; //, iinfo;
//magma_int_t lddwork = min(m, n);
//float *work, *panel;
float *dwork;
//magma_queue_t stream[2];
magma_int_t ldt=nb; // need to be an input parameter
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0)
return *info;
if (MAGMA_SUCCESS != magma_smalloc( &dwork, n*nb )) {
printf ("!!!! sorgqr_2stage magma_alloc failed for: dwork\n" );
exit(-1);
}
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code after the last block.
The first kk columns are handled by the block method.
ki is start of 2nd-to-last block. */
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
/* Set A(1:kk,kk+1:n) to zero. */
/* and A(kk+1:m, kk+1:n) = I */
magmablas_slaset( MagmaFull, kk, n-kk, c_zero, c_zero, dA(0, kk), ldda );
magmablas_slaset( MagmaFull, m-kk, n-kk, c_zero, c_one, dA(kk,kk), ldda );
}
else {
ki = 0;
kk = 0;
}
/* Allocate work space on CPU in pinned memory */
//lwork = (n+m) * nb;
//if (kk < n)
// lwork = max(lwork, n * nb + (m-kk)*(n-kk));
//if (MAGMA_SUCCESS != magma_smalloc_pinned( &work, (lwork) )) {
// *info = MAGMA_ERR_HOST_ALLOC;
// return *info;
//}
//panel = work + n * nb;
//magma_queue_create( &stream[0] );
//magma_queue_create( &stream[1] );
/* Use unblocked code for the last or only block. */
if (kk < n) {
i__1 = m - kk;
i__2 = n - kk;
i__3 = k - kk;
//magma_sgetmatrix(i__1, i__2, dA(kk, kk), ldda, panel, i__1);
//lapackf77_sorgqr(&i__1, &i__2, &i__3, panel, &i__1, &tau[kk],
// work, &lwork, &iinfo);
//
//magma_ssetmatrix(i__1, i__2, panel, i__1, dA(kk, kk), ldda);
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__1, i__2, i__3,
dA(kk, kk-nb), ldda, dT(kk-nb), ldt,
dA(kk, kk), ldda, dwork, i__2);
//magmablas_slaset(MagmaFull, kk-nb, nb, c_zero, c_zero, dA(0,kk-nb), ldda);
//magmablas_slaset(MagmaFull, m-(kk-nb), nb, c_zero, c_one, dA(kk-nb,kk-nb), ldda);
}
if (kk > 0) {
/* Use blocked code */
for (i = ki; i >= nb; i -= nb) {
ib = min(nb, k - i);
/* Send current panel to the CPU for update */
i__2 = m - i;
//magma_sgetmatrix_async( i__2, ib, dA(i,i), ldda, panel, i__2, stream[0] ); // verify
if (i + ib < n) {
/* Apply H to A(i:m,i+ib:n) from the left */
i__3 = n - i;
magmablas_slaset( MagmaFull, i, ib, c_zero, c_zero, dA(0,i), ldda );
magmablas_slaset( MagmaFull, m-i, ib, c_zero, c_one, dA(i,i), ldda );
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__2, i__3, ib,
dA(i, i-nb), ldda, dT(i-nb), ldt,
dA(i, i), ldda, dwork, i__3);
}
/* Apply H to rows i:m of current block on the CPU */
//magma_queue_sync( stream[0] );
//lapackf77_sorgqr(&i__2, &ib, &ib, panel, &i__2, &tau[i],
// work, &lwork, &iinfo);
//magma_ssetmatrix_async( i__2, ib, panel, i__2, dA(i,i), ldda, stream[1] ); // verify
/* Set rows 1:i-1 of current block to zero */
i__2 = i + ib;
//magmablas_slaset(MagmaFull, i-ib, ib, c_zero, c_zero, dA(0,i-ib), ldda);
//magmablas_slaset(MagmaFull, m-(i-ib), ib, c_zero, c_one, dA(i-ib,i-ib), ldda);
}
}
magmablas_slaset( MagmaFull, m, nb, c_zero, c_one, dA(0,0), ldda );
magma_free( dwork );
//magma_free_pinned( work );
//magma_queue_destroy( stream[0] );
//magma_queue_destroy( stream[1] );
return *info;
} /* magma_sorgqr_gpu */
/**
Purpose
-------
SGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
It uses 2 queues to overlap communication and computation.
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,out]
A REAL array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_sgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_sgetrf(
magma_int_t m, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *ipiv,
magma_int_t *info)
{
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_)*nb + (j_)*nb*ldda + dA_offset)
#define dAT(i_, j_) dAT, ((i_)*nb*lddat + (j_)*nb + dAT_offset)
#define dwork(i_) dwork, (i_)
#else
#define dA(i_, j_) ( dA + (i_)*nb + (j_)*nb*ldda)
#define dAT(i_, j_) ( dAT + (i_)*nb*lddat + (j_)*nb)
#define dwork(i_) (dwork + (i_))
#endif
// Constants
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
// Local variables
float *work;
magmaFloat_ptr dA, dAT, dwork;
magma_int_t iinfo, nb;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
nb = magma_get_sgetrf_nb( m, n );
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_sgetrf( &m, &n, A, &lda, ipiv, info );
}
else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, lddat, maxdim;
magma_int_t i, j, rows, cols, s = min(m, n)/nb;
maxm = magma_roundup( m, 32 );
maxn = magma_roundup( n, 32 );
maxdim = max( maxm, maxn );
lddat = maxn;
ldda = maxm;
/* set number of GPUs */
magma_int_t ngpu = magma_num_gpus();
if ( ngpu > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
magma_queue_t queues[2] = { NULL, NULL };
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
/* check the memory requirement */
size_t mem_size = magma_queue_mem_size( queues[0] );
mem_size /= sizeof(float);
magma_int_t h = 1+(2+ngpu);
magma_int_t ngpu2 = ngpu;
magma_int_t NB = (magma_int_t)(0.8*mem_size/maxm - h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
if ( ngpu > ceil((float)NB/nb) ) {
ngpu2 = (magma_int_t)ceil((float)NB/nb);
h = 1+(2+ngpu2);
NB = (magma_int_t)(0.8*mem_size/maxm - h*nb);
}
if ( ngpu2*NB < n ) {
/* require too much memory, so call non-GPU-resident version */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
work = A;
if (maxdim*maxdim < 2*maxm*maxn) {
// if close to square, allocate square matrix and transpose in-place
// dwork is nb*maxm for panel, and maxdim*maxdim for A
if (MAGMA_SUCCESS != magma_smalloc( &dwork, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
dA = dwork + nb*maxm;
ldda = lddat = maxdim;
magma_ssetmatrix( m, n, A, lda, dA(0,0), ldda, queues[0] );
dAT = dA;
magmablas_stranspose_inplace( maxdim, dAT(0,0), lddat, queues[0] );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
// dwork is nb*maxm for panel, and maxm*maxn for A
if (MAGMA_SUCCESS != magma_smalloc( &dwork, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
dA = dwork + nb*maxm;
magma_ssetmatrix( m, n, A, lda, dA(0,0), ldda, queues[0] );
if (MAGMA_SUCCESS != magma_smalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dwork );
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
magmablas_stranspose( m, n, dA(0,0), ldda, dAT(0,0), lddat, queues[0] );
}
lapackf77_sgetrf( &m, &nb, work, &lda, ipiv, &iinfo );
for( j = 0; j < s; j++ ) {
// get j-th panel from device
cols = maxm - j*nb;
if (j > 0) {
magmablas_stranspose( nb, cols, dAT(j,j), lddat, dwork(0), cols, queues[0] );
magma_queue_sync( queues[0] );
magma_sgetmatrix_async( m-j*nb, nb, dwork(0), cols, work, lda, queues[1] );
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (j+1)*nb, nb,
c_one, dAT(j-1,j-1), lddat,
dAT(j-1,j+1), lddat, queues[0] );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-j*nb, nb,
c_neg_one, dAT(j-1,j+1), lddat,
dAT(j, j-1), lddat,
c_one, dAT(j, j+1), lddat, queues[0] );
// do the cpu part
rows = m - j*nb;
magma_queue_sync( queues[1] );
lapackf77_sgetrf( &rows, &nb, work, &lda, ipiv+j*nb, &iinfo );
}
if (*info == 0 && iinfo > 0)
*info = iinfo + j*nb;
// put j-th panel onto device
magma_ssetmatrix_async( m-j*nb, nb, work, lda, dwork(0), cols, queues[1] );
for( i=j*nb; i < j*nb + nb; ++i ) {
ipiv[i] += j*nb;
}
magmablas_slaswp( n, dAT(0,0), lddat, j*nb + 1, j*nb + nb, ipiv, 1, queues[0] );
magma_queue_sync( queues[1] );
magmablas_stranspose( cols, nb, dwork(0), cols, dAT(j,j), lddat, queues[0] );
// do the small non-parallel computations (next panel update)
if (s > (j+1)) {
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat, queues[0] );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat, queues[0] );
}
else {
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat, queues[0] );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat, queues[0] );
}
}
magma_int_t nb0 = min( m - s*nb, n - s*nb );
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_stranspose( nb0, rows, dAT(s,s), lddat, dwork(0), cols, queues[0] );
magma_sgetmatrix_async( rows, nb0, dwork(0), cols, work, lda, queues[0] );
magma_queue_sync( queues[0] );
// do the cpu part
lapackf77_sgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo );
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
for( i=s*nb; i < s*nb + nb0; ++i ) {
ipiv[i] += s*nb;
}
magmablas_slaswp( n, dAT(0,0), lddat, s*nb + 1, s*nb + nb0, ipiv, 1, queues[0] );
// put j-th panel onto device
magma_ssetmatrix_async( rows, nb0, work, lda, dwork(0), cols, queues[0] );
magmablas_stranspose( rows, nb0, dwork(0), cols, dAT(s,s), lddat, queues[0] );
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s, s), lddat,
dAT(s, s)+nb0, lddat, queues[0] );
}
// undo transpose
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_stranspose_inplace( maxdim, dAT(0,0), lddat, queues[0] );
magma_sgetmatrix( m, n, dAT(0,0), lddat, A, lda, queues[0] );
}
else {
magmablas_stranspose( n, m, dAT(0,0), lddat, dA(0,0), ldda, queues[0] );
magma_sgetmatrix( m, n, dA(0,0), ldda, A, lda, queues[0] );
magma_free( dAT );
}
magma_free( dwork );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
}
return *info;
} /* magma_sgetrf */
BlockOperator & ParBlockNonlinearForm::GetGradient(const Vector &x) const
{
if (pBlockGrad == NULL)
{
pBlockGrad = new BlockOperator(block_trueOffsets);
}
Array<const ParFiniteElementSpace *> pfes(fes.Size());
for (int s1=0; s1<fes.Size(); ++s1)
{
pfes[s1] = ParFESpace(s1);
for (int s2=0; s2<fes.Size(); ++s2)
{
phBlockGrad(s1,s2)->Clear();
}
}
GetLocalGradient(x); // gradients are stored in 'Grads'
if (fnfi.Size() > 0)
{
MFEM_ABORT("TODO: assemble contributions from shared face terms");
}
for (int s1=0; s1<fes.Size(); ++s1)
{
for (int s2=0; s2<fes.Size(); ++s2)
{
OperatorHandle dA(phBlockGrad(s1,s2)->Type()),
Ph(phBlockGrad(s1,s2)->Type()),
Rh(phBlockGrad(s1,s2)->Type());
if (s1 == s2)
{
dA.MakeSquareBlockDiag(pfes[s1]->GetComm(), pfes[s1]->GlobalVSize(),
pfes[s1]->GetDofOffsets(), Grads(s1,s1));
Ph.ConvertFrom(pfes[s1]->Dof_TrueDof_Matrix());
phBlockGrad(s1,s1)->MakePtAP(dA, Ph);
}
else
{
dA.MakeRectangularBlockDiag(pfes[s1]->GetComm(),
pfes[s1]->GlobalVSize(),
pfes[s2]->GlobalVSize(),
pfes[s1]->GetDofOffsets(),
pfes[s2]->GetDofOffsets(),
Grads(s1,s2));
Rh.ConvertFrom(pfes[s1]->Dof_TrueDof_Matrix());
Ph.ConvertFrom(pfes[s2]->Dof_TrueDof_Matrix());
phBlockGrad(s1,s2)->MakeRAP(Rh, dA, Ph);
}
pBlockGrad->SetBlock(s1, s2, phBlockGrad(s1,s2)->Ptr());
}
}
return *pBlockGrad;
}
/**
Purpose
-------
ZHETRD reduces a complex Hermitian matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = MagmaUpper, the diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, and the elements below the first subdiagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
d COMPLEX_16 array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e COMPLEX_16 array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
tau COMPLEX_16 array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX_16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= N*NB, where NB is the
optimal blocksize given by magma_get_zhetrd_nb().
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_zheev_comp
********************************************************************/
extern "C" magma_int_t
magma_zhetrd(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda,
double *d, double *e, magmaDoubleComplex *tau,
magmaDoubleComplex *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_, j_) ( A + (i_) + (j_)*lda )
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ldda = roundup( n, 32 );
magma_int_t nb = magma_get_zhetrd_nb( n );
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const double d_one = MAGMA_D_ONE;
magma_int_t kk, nx;
magma_int_t i, j, i_n;
magma_int_t iinfo;
magma_int_t ldw, lddw, lwkopt;
magma_int_t lquery;
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
}
/* Determine the block size. */
ldw = n;
lddw = ldda;
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_Z_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaDoubleComplex *dA;
#ifdef FAST_HEMV
magma_int_t ldwork2 = ldda*ceildiv(n,64);
#else
magma_int_t ldwork2 = 0;
#endif
if (MAGMA_SUCCESS != magma_zmalloc( &dA, ldda*n + 2*lddw*nb + ldwork2 )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaDoubleComplex *dwork = dA + ldda*n;
#ifdef FAST_HEMV
magmaDoubleComplex *dwork2 = dwork + 2*lddw*nb;
#endif
//if (n < 2048)
// nx = n;
//else
// nx = 512;
nx = min( 128, n ); // nx <= n is required
// clear out dwork in case it has NANs (used as y in zhemv)
// rest of dwork (used as work in magmablas_zhemv) doesn't need to be cleared
magmablas_zlaset( MagmaFull, n, nb, c_zero, c_zero, dwork, lddw );
if (upper) {
/* Copy the matrix to the GPU */
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
kk = n - (n - nx + nb - 1) / nb * nb;
for (i = n - nb; i >= kk; i -= nb) {
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != n-nb)
magma_zgetmatrix( i+nb, nb, dA(0, i), ldda, A(0, i), lda );
#ifdef FAST_HEMV
magma_zlatrd2( uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldw, dA(0, 0), ldda, dwork, lddw,
dwork2, ldwork2 );
#else
magma_zlatrd( uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldw, dA(0, 0), ldda, dwork, lddw );
#endif
/* Update the unreduced submatrix A(0:i-2,0:i-2), using an
update of the form: A := A - V*W' - W*V' */
magma_zsetmatrix( i + nb, nb, work, ldw, dwork, lddw );
magma_zher2k( uplo, MagmaNoTrans, i, nb, c_neg_one,
dA(0, i), ldda, dwork, lddw,
d_one, dA(0, 0), ldda );
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
*A(j-1,j) = MAGMA_Z_MAKE( e[j - 1], 0 );
d[j] = MAGMA_Z_REAL( *A(j, j) );
}
}
magma_zgetmatrix( kk, kk, dA(0, 0), ldda, A(0, 0), lda );
/* Use CPU code to reduce the last or only block */
lapackf77_zhetrd( uplo_, &kk, A(0, 0), &lda, d, e, tau, work, &lwork, &iinfo );
}
else {
/* Copy the matrix to the GPU */
if (1 <= n-nx)
magma_zsetmatrix( n, n, A(0,0), lda, dA(0,0), ldda );
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb) {
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != 0)
magma_zgetmatrix( n-i, nb, dA(i, i), ldda, A(i, i), lda );
#ifdef FAST_HEMV
magma_zlatrd2( uplo, n-i, nb, A(i, i), lda, &e[i], &tau[i],
work, ldw, dA(i, i), ldda, dwork, lddw,
dwork2, ldwork2 );
#else
magma_zlatrd( uplo, n-i, nb, A(i, i), lda, &e[i], &tau[i],
work, ldw, dA(i, i), ldda, dwork, lddw );
#endif
/* Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V' */
magma_zsetmatrix( n-i, nb, work, ldw, dwork, lddw );
magma_zher2k( MagmaLower, MagmaNoTrans, n-i-nb, nb, c_neg_one,
dA(i+nb, i), ldda, &dwork[nb], lddw,
d_one, dA(i+nb, i+nb), ldda );
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
*A(j+1,j) = MAGMA_Z_MAKE( e[j], 0 );
d[j] = MAGMA_Z_REAL( *A(j, j) );
}
}
/* Use CPU code to reduce the last or only block */
if (1 <= n-nx)
magma_zgetmatrix( n-i, n-i, dA(i, i), ldda, A(i, i), lda );
i_n = n-i;
lapackf77_zhetrd( uplo_, &i_n, A(i, i), &lda, &d[i], &e[i],
&tau[i], work, &lwork, &iinfo );
}
magma_free( dA );
work[0] = MAGMA_Z_MAKE( lwkopt, 0 );
return *info;
} /* magma_zhetrd */
extern "C" magma_int_t
magma_zhegst(magma_int_t itype, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *b, magma_int_t ldb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZHEGST reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U**H or L**H*A*L.
B must have been previously factorized as U**H*U or L*L**H by ZPOTRF.
Arguments
=========
ITYPE (input) INTEGER
= 1: compute inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H);
= 2 or 3: compute U*A*U**H or L**H*A*L.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored and B is factored as
U**H*U;
= 'L': Lower triangle of A is stored and B is factored as
L*L**H.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input) COMPLEX_16 array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
=====================================================================*/
char uplo_[2] = {uplo, 0};
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_half = MAGMA_Z_HALF;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex *dw;
magma_int_t ldda = n;
magma_int_t lddb = n;
double d_one = 1.0;
int upper = lapackf77_lsame(uplo_, "U");
/* Test the input parameters. */
*info = 0;
if (itype<1 || itype>3){
*info = -1;
}else if ((! upper) && (! lapackf77_lsame(uplo_, "L"))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
}else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
if (MAGMA_SUCCESS != magma_zmalloc( &dw, 2*n*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
nb = magma_get_zhegst_nb(n);
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
magma_zsetmatrix( n, n, B(0, 0), ldb, dB(0, 0), lddb );
/* Use hybrid blocked code */
if (itype==1) {
if (upper) {
/* Compute inv(U')*A*inv(U) */
for(k = 0; k<n; k+=nb){
kb = min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_zher2k(MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_ztrsm(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one ,dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
/* Compute inv(L)*A*inv(L') */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_zher2k(MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda);
}
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one ,dB(0,0), lddb,
dA(0,k), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ztrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k, k), &lda, B(k, k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
else {
/* Compute L'*A*L */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one ,dB(0,0), lddb,
dA(k,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ztrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_zgetmatrix( n, n, dA(0, 0), ldda, A(0, 0), lda );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dw );
return *info;
} /* magma_zhegst_gpu */
/**
Purpose
=======
SSYTRF_nopiv computes the LDLt factorization of a real symmetric
matrix A. This version does not require work space on the GPU passed
as input. GPU memory is allocated in the routine.
The factorization has the form
A = U^H * D * U, if UPLO = MagmaUpper, or
A = L * D * L^H, if UPLO = MagmaLower,
where U is an upper triangular matrix, L is lower triangular, and
D is a diagonal matrix.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U^H D U or A = L D L^H.
\n
Higher performance is achieved if A is in pinned memory.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -6, the GPU memory allocation failed
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_ssysv_comp
******************************************************************* */
extern "C" magma_int_t
magma_ssytrf_nopiv(
magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i, j) ( A +(j)*lda + (i))
#define dA(i, j) (dA +(j)*ldda + (i))
#define dW(i, j) (dW +(j)*ldda + (i))
#define dWt(i, j) (dW +(j)*nb + (i))
/* Constants */
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
/* Local variables */
bool upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, ldda, nb, ib, iinfo;
magmaFloat_ptr dA;
magmaFloat_ptr dW;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
/* Quick return */
if ( n == 0 )
return MAGMA_SUCCESS;
ldda = magma_roundup( n, 32 );
nb = magma_get_ssytrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
if ((MAGMA_SUCCESS != magma_smalloc(&dA, n *ldda)) ||
(MAGMA_SUCCESS != magma_smalloc(&dW, nb*ldda))) {
/* alloc failed so call the non-GPU-resident version */
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_device_t cdev;
magma_queue_t queues[2];
magma_event_t event;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magma_event_create( &event );
trace_init( 1, 1, 2, queues );
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// copy matrix to GPU
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(j+jb, jb, A(0, j), lda, dA(0, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
}
// main loop
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
if ( j != 0) {
//magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, queues[1]);
}
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(queues[1]);
trace_cpu_start( 0, "potrf", "potrf" );
magma_ssytrf_nopiv_cpu( MagmaUpper, jb, ib, A(j, j), lda, info );
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
// copy j-th column of U back to CPU
trace_gpu_start( 0, 1, "get", "get" );
magma_sgetmatrix_async(j, jb, dA(0, j), ldda, A(0, j), lda, queues[1]);
trace_gpu_end( 0, 1 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaUnit,
jb, (n-j-jb),
c_one, dA(j, j), ldda,
dA(j, j+jb), ldda,
queues[0] );
magma_scopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb, queues[0] );
// update the trailing submatrix with D
magmablas_slascl_diag( MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
queues[0], &iinfo);
trace_gpu_end( 0, 0 );
// update the trailing submatrix with U and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k < n; k += nb) {
magma_int_t kb = min(nb,n-k);
magma_sgemm( MagmaConjTrans, MagmaNoTrans, kb, n-k, jb,
c_neg_one, dWt(0, k), nb,
dA(j, k), ldda,
c_one, dA(k, k), ldda,
queues[0]);
if (k == j+jb) {
// magma_event_record( event, queues[0] );
magma_queue_sync( queues[0] );
}
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// copy the matrix to GPU
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async((n-j), jb, A(j, j), lda, dA(j, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
}
// main loop
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
if (j != 0) {
//magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, queues[1]);
}
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(queues[1]);
trace_cpu_start( 0, "potrf", "potrf" );
magma_ssytrf_nopiv_cpu( MagmaLower, jb, ib, A(j, j), lda, info );
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, queues[0]);
trace_gpu_end( 0, 0 );
// copy j-th row of L back to CPU
trace_gpu_start( 0, 1, "get", "get" );
magma_sgetmatrix_async(jb, j, dA(j, 0), ldda, A(j, 0), lda, queues[1]);
trace_gpu_end( 0, 1 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda,
queues[0] );
magma_scopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda, queues[0] );
// update the trailing submatrix with D
magmablas_slascl_diag( MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
queues[0], &iinfo );
trace_gpu_end( 0, 0 );
// update the trailing submatrix with L and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k < n; k += nb) {
magma_int_t kb = min(nb,n-k);
magma_sgemm( MagmaNoTrans, MagmaConjTrans, n-k, kb, jb,
c_neg_one, dA(k, j), ldda,
dW(k, 0), ldda,
c_one, dA(k, k), ldda,
queues[0] );
if (k == j+jb) {
//magma_event_record( event, queues[0] );
magma_queue_sync(queues[0]);
}
}
trace_gpu_end( 0, 0 );
}
}
}
trace_finalize( "ssytrf.svg","trace.css" );
magma_queue_destroy(queues[0]);
magma_queue_destroy(queues[1]);
magma_event_destroy( event );
magma_free(dW);
magma_free(dA);
return MAGMA_SUCCESS;
} /* magma_ssytrf_nopiv */
/**
Purpose
-------
SGEBRD reduces a general real M-by-N matrix A to upper or lower
bidiagonal form B by an orthogonal transformation: Q**H * A * P = B.
If m >= n, B is upper bidiagonal; if m < n, B is lower bidiagonal.
Arguments
---------
@param[in]
m INTEGER
The number of rows in the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns in the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the M-by-N general matrix to be reduced.
On exit,
if m >= n, the diagonal and the first superdiagonal are
overwritten with the upper bidiagonal matrix B; the
elements below the diagonal, with the array TAUQ, represent
the orthogonal matrix Q as a product of elementary
reflectors, and the elements above the first superdiagonal,
with the array TAUP, represent the orthogonal matrix P as
a product of elementary reflectors;
\n
if m < n, the diagonal and the first subdiagonal are
overwritten with the lower bidiagonal matrix B; the
elements below the first subdiagonal, with the array TAUQ,
represent the orthogonal matrix Q as a product of
elementary reflectors, and the elements above the diagonal,
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[out]
d real array, dimension (min(M,N))
The diagonal elements of the bidiagonal matrix B:
D(i) = A(i,i).
@param[out]
e real array, dimension (min(M,N)-1)
The off-diagonal elements of the bidiagonal matrix B:
if m >= n, E(i) = A(i,i+1) for i = 1,2,...,n-1;
if m < n, E(i) = A(i+1,i) for i = 1,2,...,m-1.
@param[out]
tauq REAL array dimension (min(M,N))
The scalar factors of the elementary reflectors which
represent the orthogonal matrix Q. See Further Details.
@param[out]
taup REAL array, dimension (min(M,N))
The scalar factors of the elementary reflectors which
represent the orthogonal matrix P. See Further Details.
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= (M+N)*NB, where NB
is the optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
---------------
The matrices Q and P are represented as products of elementary
reflectors:
If m >= n,
Q = H(1) H(2) . . . H(n) and P = G(1) G(2) . . . G(n-1)
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;
v(1:i-1) = 0, v(i) = 1, and v(i+1:m) is stored on exit in A(i+1:m,i);
u(1:i) = 0, u(i+1) = 1, and u(i+2:n) is stored on exit in A(i,i+2:n);
tauq is stored in TAUQ(i) and taup in TAUP(i).
If m < n,
Q = H(1) H(2) . . . H(m-1) and P = G(1) G(2) . . . G(m)
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;
v(1:i) = 0, v(i+1) = 1, and v(i+2:m) is stored on exit in A(i+2:m,i);
u(1:i-1) = 0, u(i) = 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).
The contents of A on exit are illustrated by the following examples:
@verbatim
m = 6 and n = 5 (m > n): m = 5 and n = 6 (m < n):
( d e u1 u1 u1 ) ( d u1 u1 u1 u1 u1 )
( v1 d e u2 u2 ) ( e d u2 u2 u2 u2 )
( v1 v2 d e u3 ) ( v1 e d u3 u3 u3 )
( v1 v2 v3 d e ) ( v1 v2 e d u4 u4 )
( v1 v2 v3 v4 d ) ( v1 v2 v3 e d u5 )
( v1 v2 v3 v4 v5 )
@endverbatim
where d and e denote diagonal and off-diagonal elements of B, vi
denotes an element of the vector defining H(i), and ui an element of
the vector defining G(i).
@ingroup magma_sgesvd_comp
********************************************************************/
extern "C" magma_int_t
magma_sgebrd(
magma_int_t m, magma_int_t n,
float *A, magma_int_t lda, float *d, float *e,
float *tauq, float *taup,
float *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float *dA, *dwork;
magma_int_t ncol, nrow, jmax, nb, ldda;
magma_int_t i, j, nx;
magma_int_t iinfo;
magma_int_t minmn;
magma_int_t ldwrkx, ldwrky, lwkopt;
magma_int_t lquery;
nb = magma_get_sgebrd_nb(n);
ldda = m;
lwkopt = (m + n) * nb;
work[0] = MAGMA_S_MAKE( lwkopt, 0. );
lquery = (lwork == -1);
/* Check arguments */
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < lwkopt && (! lquery) ) {
*info = -10;
}
if (*info < 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
minmn = min(m,n);
if (minmn == 0) {
work[0] = c_one;
return *info;
}
if (MAGMA_SUCCESS != magma_smalloc( &dA, n*ldda + (m + n)*nb )) {
fprintf (stderr, "!!!! device memory allocation error in sgebrd\n" );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork = dA + n*ldda;
ldwrkx = m;
ldwrky = n;
/* Set the block/unblock crossover point NX. */
nx = 128;
/* Copy the matrix to the GPU */
if (minmn - nx >= 1) {
magma_ssetmatrix( m, n, A, lda, dA, ldda );
}
for (i=0; i < (minmn - nx); i += nb) {
/* Reduce rows and columns i:i+nb-1 to bidiagonal form and return
the matrices X and Y which are needed to update the unreduced
part of the matrix */
nrow = m - i;
ncol = n - i;
/* Get the current panel (no need for the 1st iteration) */
if ( i > 0 ) {
magma_sgetmatrix( nrow, nb, dA(i, i), ldda, A( i, i), lda );
magma_sgetmatrix( nb, ncol - nb,
dA(i, i+nb), ldda,
A( i, i+nb), lda );
}
magma_slabrd_gpu(nrow, ncol, nb,
A(i, i), lda, dA(i, i), ldda,
d+i, e+i, tauq+i, taup+i,
work, ldwrkx, dwork, ldwrkx, // x, dx
work+(ldwrkx*nb), ldwrky, dwork+(ldwrkx*nb), ldwrky); // y, dy
/* Update the trailing submatrix A(i+nb:m,i+nb:n), using an update
of the form A := A - V*Y' - X*U' */
nrow = m - i - nb;
ncol = n - i - nb;
// Send Y back to the GPU
magma_ssetmatrix( nrow, nb, work + nb, ldwrkx, dwork + nb, ldwrkx );
magma_ssetmatrix( ncol, nb,
work + (ldwrkx+1)*nb, ldwrky,
dwork + (ldwrkx+1)*nb, ldwrky );
magma_sgemm( MagmaNoTrans, MagmaConjTrans,
nrow, ncol, nb,
c_neg_one, dA(i+nb, i ), ldda,
dwork+(ldwrkx+1)*nb, ldwrky,
c_one, dA(i+nb, i+nb), ldda);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
nrow, ncol, nb,
c_neg_one, dwork+nb, ldwrkx,
dA( i, i+nb ), ldda,
c_one, dA( i+nb, i+nb ), ldda);
/* Copy diagonal and off-diagonal elements of B back into A */
if (m >= n) {
jmax = i + nb;
for (j = i; j < jmax; ++j) {
*A(j, j ) = MAGMA_S_MAKE( d[j], 0. );
*A(j, j+1) = MAGMA_S_MAKE( e[j], 0. );
}
} else {
jmax = i + nb;
for (j = i; j < jmax; ++j) {
*A(j, j ) = MAGMA_S_MAKE( d[j], 0. );
*A(j+1, j ) = MAGMA_S_MAKE( e[j], 0. );
}
}
}
/* Use unblocked code to reduce the remainder of the matrix */
nrow = m - i;
ncol = n - i;
if ( 0 < minmn - nx ) {
magma_sgetmatrix( nrow, ncol, dA(i, i), ldda, A(i, i), lda );
}
lapackf77_sgebrd( &nrow, &ncol,
A(i, i), &lda, d+i, e+i,
tauq+i, taup+i, work, &lwork, &iinfo);
work[0] = MAGMA_S_MAKE( lwkopt, 0. );
magma_free( dA );
return *info;
} /* magma_sgebrd */
extern "C" magma_err_t
magma_cgeqrf(magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info,
magma_queue_t* queue )
{
/* -- clMAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
CGEQRF computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
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.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= max( N*NB, 2*NB*NB ),
where NB can be obtained through magma_get_cgeqrf_nb(M).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) dA, dA_offset + (i) + (j)*ldda
magmaFloatComplex_ptr dA, dwork, dT;
size_t dA_offset, dwork_offset, dT_offset;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t i, k, lddwork, old_i, old_ib;
magma_int_t ib, ldda;
*info = 0;
magma_int_t nb = magma_get_cgeqrf_nb(min(m, n));
// need 2*nb*nb to store T and upper triangle of V simultaneously
magma_int_t lwkopt = max(n*nb, 2*nb*nb);
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1, lwkopt) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
// largest N for larfb is n-nb (trailing matrix lacks 1st panel)
lddwork = ((n+31)/32)*32 - nb;
ldda = ((m+31)/32)*32;
magma_int_t num_gpus = magma_num_gpus();
if( num_gpus > 1 ) {
/* call multiple-GPU interface */
printf("multiple-GPU verison not implemented\n");
return MAGMA_ERR_NOT_IMPLEMENTED;
//return magma_cgeqrf4(num_gpus, m, n, A, lda, tau, work, lwork, info);
}
// allocate space for dA, dwork, and dT
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n*ldda + nb*lddwork + nb*nb) )) {
/* Switch to the "out-of-core" (out of GPU-memory) version */
printf("non-GPU-resident version not implemented\n");
return MAGMA_ERR_NOT_IMPLEMENTED;
//return magma_cgeqrf_ooc(m, n, A, lda, tau, work, lwork, info);
}
dA_offset = 0;
dwork = dA;
dwork_offset = n*ldda;
dT = dA;
dT_offset = n*ldda + nb*lddwork;
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
Asynchronously send the matrix to the GPU except the first panel. */
magma_csetmatrix_async( m, n-nb,
A(0,nb), 0, lda,
dA(0,nb), ldda, queue[0], NULL );
old_i = 0;
old_ib = nb;
for (i = 0; i < k-nb; i += nb) {
ib = min(k-i, nb);
if (i>0) {
/* download i-th panel */
magma_queue_sync( queue[1] );
magma_cgetmatrix_async( m-i, ib,
dA(i,i), ldda,
A(i,i), 0, lda, queue[0], NULL );
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i), ldda, dT, dT_offset, nb,
dA(old_i, old_i+2*old_ib), ldda, dwork, dwork_offset, lddwork, queue[1]);
magma_cgetmatrix_async( i, ib,
dA(0,i), ldda,
A(0,i), 0, lda, queue[1], NULL );
magma_queue_sync( queue[0] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(i,i), &lda, tau+i, work, &ib);
cpanel_to_q(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
/* download the i-th V matrix */
magma_csetmatrix_async( rows, ib, A(i,i), 0, lda, dA(i,i), ldda, queue[0], NULL );
/* download the T matrix */
magma_queue_sync( queue[1] );
magma_csetmatrix_async( ib, ib, work, 0, ib, dT, dT_offset, nb, queue[0], NULL );
magma_queue_sync( queue[0] );
if (i + ib < n) {
if (i+ib < k-nb) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left (look-ahead) */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT, dT_offset, nb,
dA(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue[1]);
cq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
else {
/* After last panel, update whole trailing matrix. */
/* Apply H' to A(i:m,i+ib:n) from the left */
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dT, dT_offset, nb,
dA(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue[1]);
cq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
if (i != 0) {
magma_cgetmatrix( m, ib, dA(0,i), ldda, A(0,i), 0, lda, queue[1] );
}
magma_int_t rows = m-i;
lapackf77_cgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
}
magma_queue_sync(queue[0]);
magma_queue_sync(queue[1]);
magma_free( dA );
return *info;
} /* magma_cgeqrf */
/**
Purpose
-------
DORGQR generates an M-by-N DOUBLE_PRECISION matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF.
This version recomputes the T matrices on the CPU and sends them to the GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
A DOUBLE_PRECISION array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by DGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
lda INTEGER
The first dimension of the array A. LDA >= max(1,M).
@param[in]
tau DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF_GPU.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument has an illegal value
@ingroup magma_dgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_dorgqr2(magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
magma_int_t nb = magma_get_dgeqrf_nb(min(m, n));
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk; //, iinfo;
magma_int_t lddwork;
double *dA, *dV, *dW, *dT, *T;
double *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate GPU work space
// ldda*n for matrix dA
// ldda*nb for dV
// lddwork*nb for dW larfb workspace
ldda = ((m + 31) / 32) * 32;
lddwork = ((n + 31) / 32) * 32;
if (MAGMA_SUCCESS != magma_dmalloc( &dA, ldda*n + ldda*nb + lddwork*nb + nb*nb)) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dA + ldda*n;
dW = dA + ldda*n + ldda*nb;
dT = dA + ldda*n + ldda*nb + lddwork*nb;
// Allocate CPU work space
lwork = (n+m+nb) * nb;
magma_dmalloc_cpu( &work, lwork );
T = work;
if (work == NULL) {
magma_free( dA );
magma_free_cpu( work );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
double *V = work + (n+nb)*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
/*
lapackf77_dorgqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_dlacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_dlaset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_dlarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
V, &m_kk, work, &k_kk, A(kk, kk), &lda, work+k_kk*k_kk, &n_kk );
if (kk > 0) {
magma_dsetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_dlaset( MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda );
}
}
if (kk > 0) {
// Use blocked code
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min(nb, k - i);
// Send current panel to the GPU
mi = m - i;
lapackf77_dlaset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_dsetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&mi, &ib,
A(i,i), &lda, &tau[i], T, &nb);
magma_dsetmatrix_async( ib, ib,
T, nb,
dT, nb, stream );
// set panel to identity
magmablas_dlaset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda );
magmablas_dlaset( MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda );
magma_queue_sync( stream );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT, nb,
dA(i, i), ldda, dW, lddwork );
}
}
// copy result back to CPU
magma_dgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda);
}
magmablasSetKernelStream( NULL );
magma_queue_destroy( stream );
magma_free( dA );
magma_free_cpu( work );
return *info;
} /* magma_dorgqr */
extern "C" magma_int_t
magma_cungqr(magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *dT, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CUNGQR generates an M-by-N COMPLEX matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by CGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) COMPLEX array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by CGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) COMPLEX array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by CGEQRF_GPU.
DT (input) COMPLEX array on the GPU device.
DT contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_cgeqrf_gpu.
NB (input) INTEGER
This is the block size used in CGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk; //, iinfo;
magma_int_t lddwork;
magmaFloatComplex *dA, *dV, *dW;
magmaFloatComplex *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate GPU work space
// ldda*n for matrix dA
// ldda*nb for dV
// lddwork*nb for dW larfb workspace
ldda = ((m + 31) / 32) * 32;
lddwork = ((n + 31) / 32) * 32;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, ldda*n + ldda*nb + lddwork*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dA + ldda*n;
dW = dA + ldda*n + ldda*nb;
// Allocate CPU work space
lwork = (n+m+nb) * nb;
magma_cmalloc_cpu( &work, lwork );
if (work == NULL) {
magma_free( dA );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmaFloatComplex *V = work + (n+nb)*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
/*
// Replacing this with the following 4 routines works but cungqr is slow for
// k smaller than the cungqr's blocking size (new version can be up to 60x faster)
lapackf77_cungqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_clacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_claset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_clarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
V, &m_kk, work, &k_kk, A(kk, kk), &lda, work+k_kk*k_kk, &n_kk );
if (kk > 0) {
magma_csetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_claset( MagmaUpperLower, kk, n - kk, dA(0, kk), ldda );
}
}
if (kk > 0) {
// Use blocked code
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min(nb, k - i);
// Send current panel to the GPU
mi = m - i;
lapackf77_claset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_csetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
// set panel to identity
magmablas_claset( MagmaUpperLower, i, ib, dA(0, i), ldda );
magmablas_claset_identity( mi, ib, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_clarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
// copy result back to CPU
magma_cgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda);
}
magmablasSetKernelStream( NULL );
magma_queue_destroy( stream );
magma_free( dA );
magma_free_cpu( work );
return *info;
} /* magma_cungqr */
/**
Purpose
-------
SLAHRU is an auxiliary MAGMA routine that is used in SGEHRD to update
the trailing sub-matrices after the reductions of the corresponding
panels.
See further details below.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ihi INTEGER
Last row to update. Same as IHI in sgehrd.
@param[in]
k INTEGER
Number of rows of the matrix Am (see details below)
@param[in]
nb INTEGER
Block size
@param[out]
A REAL array, dimension (LDA,N-K)
On entry, the N-by-(N-K) general matrix to be updated. The
computation is done on the GPU. After Am is updated on the GPU
only Am(1:NB) is transferred to the CPU - to update the
corresponding Am matrix. See Further Details below.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
dA REAL array on the GPU, dimension (LDDA,N-K).
On entry, the N-by-(N-K) general matrix to be updated.
On exit, the 1st K rows (matrix Am) of A are updated by
applying an orthogonal transformation from the right
Am = Am (I-V T V'), and sub-matrix Ag is updated by
Ag = (I - V T V') Ag (I - V T V(NB+1:)' )
where Q = I - V T V' represent the orthogonal matrix
(as a product of elementary reflectors V) used to reduce
the current panel of A to upper Hessenberg form. After Am
is updated Am(:,1:NB) is sent to the CPU.
See Further Details below.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[in,out]
dY (workspace) REAL array on the GPU, dimension (LDDY, NB).
On entry the (N-K)-by-NB Y = A V. It is used internally
as workspace, so its value is changed on exit.
@param[in]
lddy INTEGER
The leading dimension of the array dY. LDDY >= max(1,N).
@param[in,out]
dV (workspace) REAL array on the GPU, dimension (LDDV, NB).
On entry the (N-K)-by-NB matrix V of elementary reflectors
used to reduce the current panel of A to upper Hessenberg form.
The rest K-by-NB part is used as workspace. V is unchanged on
exit.
@param[in]
lddv INTEGER
The leading dimension of the array dV. LDDV >= max(1,N).
@param[in]
dT REAL array on the GPU, dimension (NB, NB).
On entry the NB-by-NB upper trinagular matrix defining the
orthogonal Hessenberg reduction transformation matrix for
the current panel. The lower triangular part are 0s.
@param
dwork (workspace) REAL array on the GPU, dimension N*NB.
Further Details
---------------
This implementation follows the 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.
The difference is that here Am is computed on the GPU.
M is renamed Am, G is renamed Ag.
@ingroup magma_sgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_slahru(
magma_int_t n, magma_int_t ihi, magma_int_t k, magma_int_t nb,
float *A, magma_int_t lda,
float *dA, magma_int_t ldda,
float *dY, magma_int_t lddy,
float *dV, magma_int_t lddv,
float *dT,
float *dwork )
{
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float *dYm = dV + ihi - k;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (ihi < 0 || ihi > n) {
info = -2;
} else if (k < 0 || k > n) {
info = -3;
} else if (nb < 1 || nb > n) {
info = -4;
} else if (lda < max(1,n)) {
info = -6;
} else if (ldda < max(1,n)) {
info = -8;
} else if (lddy < max(1,n)) {
info = -10;
} else if (lddv < max(1,n)) {
info = -12;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
// top part of Y, above panel, hasn't been computed yet, so do that now
// Ym = Am V = A(0:k-1, 0:ihi-k-1) * V(0:ihi-k-1, 0:nb-1)
magma_sgemm( MagmaNoTrans, MagmaNoTrans, k, nb, ihi-k,
c_one, dA, ldda,
dV, lddv,
c_zero, dYm, ldda );
// -----
// on right, A := A Q = A - A V T V'
// Update Am = Am - Am V T V' = Am - Ym W', with W = V T'
// W = V T' = V(0:ihi-k-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
magma_sgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, nb, nb,
c_one, dV, lddv,
dT, nb,
c_zero, dwork, ldda );
// Am = Am - Ym W' = A(0:k-1, 0:ihi-k-1) - Ym(0:k-1, 0:nb-1) * W(0:ihi-k-1, 0:nb-1)'
magma_sgemm( MagmaNoTrans, MagmaConjTrans, k, ihi-k, nb,
c_neg_one, dYm, ldda,
dwork, ldda,
c_one, dA, ldda );
// copy first nb columns of Am, A(0:k-1, 0:nb-1), to host
magma_sgetmatrix( k, nb, dA, ldda, A, lda );
// -----
// on right, A := A Q = A - A V T V'
// Update Ag = Ag - Ag V T V' = Ag - Y W'
// Ag = Ag - Y W' = A(k:ihi-1, nb:ihi-k-1) - Y(0:ihi-k-1, 0:nb-1) * W(nb:ihi-k-1, 0:nb-1)'
magma_sgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, ihi-k-nb, nb,
c_neg_one, dY, ldda,
dwork + nb, ldda,
c_one, dA(k,nb), ldda );
// -----
// on left, A := Q' A = A - V T' V' A
// Ag2 = Ag2 - V T' V' Ag2 = W Yg, with W = V T' and Yg = V' Ag2
// Note that Ag is A(k:ihi, nb+1:ihi-k)
// while Ag2 is A(k:ihi, nb+1: n -k)
// Z = V(0:ihi-k-1, 0:nb-1)' * A(k:ihi-1, nb:n-k-1); Z is stored over Y
magma_sgemm( MagmaConjTrans, MagmaNoTrans, nb, n-k-nb, ihi-k,
c_one, dV, lddv,
dA(k,nb), ldda,
c_zero, dY, nb );
// Ag2 = Ag2 - W Z = A(k:ihi-1, nb:n-k-1) - W(nb:n-k-1, 0:nb-1) * Z(0:nb-1, nb:n-k-1)
magma_sgemm( MagmaNoTrans, MagmaNoTrans, ihi-k, n-k-nb, nb,
c_neg_one, dwork, ldda,
dY, nb,
c_one, dA(k,nb), ldda );
return info;
}
extern "C" magma_int_t
magma_dsytrf_nopiv(magma_uplo_t uplo, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *info)
{
/* -- MAGMA (version 1.6.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2011
Purpose
=======
DSYTRF_nopiv computes the LDLt factorization of a real symmetric
matrix A. This version does not require work space on the GPU passed
as input. GPU memory is allocated in the routine.
The factorization has the form
A = U\*\*H * D * U, if UPLO = 'U', or
A = L * D * L\*\*H, if UPLO = 'L',
where U is an upper triangular matrix, L is lower triangular, and
D is a diagonal matrix.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U\*\*H*U or A = L*L\*\*H.
Higher performance is achieved if A is in pinned memory, e.g.
allocated using cudaMallocHost.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -6, the GPU memory allocation failed
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
/* Local variables */
double zone = MAGMA_D_ONE;
double mzone = MAGMA_D_NEG_ONE;
int upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, ldda, nb, ib, iinfo;
magmaDouble_ptr dA;
magmaDouble_ptr dW;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
/* Quick return */
if ( n == 0 )
return MAGMA_SUCCESS;
ldda = ((n+31)/32)*32;
nb = magma_get_dsytrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
if ((MAGMA_SUCCESS != magma_dmalloc(&dA, n *ldda)) ||
(MAGMA_SUCCESS != magma_dmalloc(&dW, nb*ldda))) {
/* alloc failed so call the non-GPU-resident version */
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_event_t event;
magma_queue_create(&stream[0]);
magma_queue_create(&stream[1]);
magma_event_create( &event );
trace_init( 1, 1, 2, (CUstream_st**)stream );
//if (nb <= 1 || nb >= n)
//{
// lapackf77_dpotrf(uplo_, &n, a, &lda, info);
//} else
{
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// copy matrix to GPU
for (j=0; j<n; j+=nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_dsetmatrix_async(j+jb, jb, A(0, j), lda, dA(0, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
}
// main loop
for (j=0; j<n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
magma_dgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, stream[0]);
trace_gpu_end( 0, 0 );
// copy j-th column of U back to CPU
magma_queue_wait_event( stream[1], event );
trace_gpu_start( 0, 1, "get", "get" );
magma_dgetmatrix_async(j, jb, dA(0, j), ldda, A(0, j), lda, stream[1]);
trace_gpu_end( 0, 1 );
// factorize the diagonal block
magma_queue_sync(stream[0]);
trace_cpu_start( 0, "potrf", "potrf" );
dsytrf_nopiv_cpu(MagmaUpper, jb, ib, A(j, j), lda, info);
trace_cpu_end( 0 );
if (*info != 0){
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_dsetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_dtrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaUnit,
jb, (n-j-jb),
zone, dA(j, j), ldda,
dA(j, j+jb), ldda);
magma_dcopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb );
// update the trailing submatrix with D
magmablas_dlascl_diag(MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
&iinfo);
magma_event_record( event, stream[0] );
trace_gpu_end( 0, 0 );
// update the trailing submatrix with U and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k<n; k+=nb)
{
magma_int_t kb = min(nb,n-k);
magma_dgemm(MagmaConjTrans, MagmaNoTrans, kb, n-k, jb,
mzone, dWt(0, k), nb,
dA(j, k), ldda,
zone, dA(k, k), ldda);
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// copy the matrix to GPU
for (j=0; j<n; j+=nb) {
jb = min(nb, (n-j));
trace_gpu_start( 0, 0, "set", "set" );
magma_dsetmatrix_async((n-j), jb, A(j, j), lda, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
}
// main loop
for (j=0; j<n; j+=nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
magma_dgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), lda, stream[0]);
trace_gpu_end( 0, 0 );
// copy j-th row of L back to CPU
magma_queue_wait_event( stream[1], event );
trace_gpu_start( 0, 1, "get", "get" );
magma_dgetmatrix_async(jb, j, dA(j, 0), ldda, A(j, 0), lda, stream[1]);
trace_gpu_end( 0, 1 );
// factorize the diagonal block
magma_queue_sync(stream[0]);
trace_cpu_start( 0, "potrf", "potrf" );
dsytrf_nopiv_cpu(MagmaLower, jb, ib, A(j, j), lda, info);
trace_cpu_end( 0 );
if (*info != 0){
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_dsetmatrix_async(jb, jb, A(j, j), lda, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_dtrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaUnit,
(n-j-jb), jb,
zone, dA(j, j), ldda,
dA(j+jb, j), ldda);
magma_dcopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda );
// update the trailing submatrix with D
magmablas_dlascl_diag(MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
&iinfo);
magma_event_record( event, stream[0] );
trace_gpu_end( 0, 0 );
// update the trailing submatrix with L and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k<n; k+=nb)
{
magma_int_t kb = min(nb,n-k);
magma_dgemm(MagmaNoTrans, MagmaConjTrans, n-k, kb, jb,
mzone, dA(k, j), ldda,
dW(k, 0), ldda,
zone, dA(k, k), ldda);
}
trace_gpu_end( 0, 0 );
}
}
}
}
trace_finalize( "dsytrf.svg","trace.css" );
magma_queue_destroy(stream[0]);
magma_queue_destroy(stream[1]);
magma_event_destroy( event );
magma_free(dW);
magma_free(dA);
return MAGMA_SUCCESS;
} /* magma_dsytrf_nopiv */
extern "C" magma_err_t
magma_ssytrd(char uplo, magma_int_t n,
float *a, magma_int_t lda,
float *d, float *e, float *tau,
float *work, magma_int_t lwork,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMAGMA (version 1.0.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
April 2012
Purpose
=======
SSYTRD reduces a real symmetric matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**T * A * Q = T.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = 'U', the diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, and the elements below the first subdiagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
D (output) REAL array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
E (output) REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = 'U', E(i) = A(i+1,i) if UPLO = 'L'.
TAU (output) REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
char uplo_[2] = {uplo, 0};
magma_int_t ldda = lda;
magma_int_t nb = magma_get_ssytrd_nb(n);
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float d_one = MAGMA_D_ONE;
magma_int_t kk, nx;
magma_int_t i, j, i_n;
magma_int_t iinfo;
magma_int_t ldwork, lddwork, lwkopt;
magma_int_t lquery;
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
}
if (*info == 0) {
/* Determine the block size. */
ldwork = lddwork = n;
lwkopt = n * nb;
// ACD
// MAGMA_S_SET2REAL( work[0], lwkopt );
MAGMA_S_SET2REAL( work[0], (float) lwkopt );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaFloat_ptr da;
size_t da_offset = 0;
if (MAGMA_SUCCESS != magma_malloc( &da, (n*ldda + 2*n*nb )*sizeof(float))) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaFloat_ptr dwork = da;
size_t dwork_offset = da_offset + (n)*ldda;
if (n < 2048)
nx = n;
else
nx = 512;
if (upper) {
/* Copy the matrix to the GPU */
magma_ssetmatrix( n, n, A(0, 0), 0, lda, dA(0, 0), ldda, queue );
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
kk = n - (n - nx + nb - 1) / nb * nb;
for (i = n - nb; i >= kk; i -= nb)
{
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i!=n-nb)
magma_sgetmatrix( i+nb, nb, dA(0, i), ldda, A(0, i), 0, lda, queue );
magma_slatrd(uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldwork, dA(0, 0), ldda, dwork, dwork_offset, lddwork, queue);
/* Update the unreduced submatrix A(0:i-2,0:i-2), using an
update of the form: A := A - V*W' - W*V' */
magma_ssetmatrix( i + nb, nb, work, 0, ldwork, dwork, dwork_offset, lddwork, queue );
magma_ssyr2k(magma_uplo_const(uplo), MagmaNoTrans, i, nb, c_neg_one,
dA(0, i), ldda, dwork, dwork_offset,
lddwork, d_one, dA(0, 0), ldda, queue);
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
MAGMA_S_SET2REAL( *A(j-1, j), e[j - 1] );
d[j] = MAGMA_S_REAL( *A(j, j) );
}
}
magma_sgetmatrix( kk, kk, dA(0, 0), ldda, A(0, 0), 0, lda, queue );
/* Use unblocked code to reduce the last or only block */
lapackf77_ssytd2(uplo_, &kk, A(0, 0), &lda, d, e, tau, &iinfo);
}
else
{
/* Copy the matrix to the GPU */
if (1<=n-nx)
magma_ssetmatrix( n, n, A(0,0), 0, lda, dA(0,0), ldda, queue );
#ifdef FAST_SYMV
// TODO this leaks memory from da, above
magmaFloat_ptr dwork2;
if (MAGMA_SUCCESS != magma_malloc( &dwork2, (n*n)*sizeof(float) )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
size_t dwork2_offset = 0;
#endif
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb)
{
/* Reduce columns i:i+nb-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i!=0)
magma_sgetmatrix( n-i, nb, dA(i, i), ldda, A(i, i), 0, lda, queue );
#ifdef FAST_SYMV
// unported
magma_slatrd2(uplo, n-i, nb, A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA(i, i), ldda,
dwork, lddwork, dwork2, n*n);
#else
magma_slatrd(uplo, n-i, nb, A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA(i, i), ldda,
dwork, dwork_offset, lddwork, queue);
#endif
/* Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V' */
magma_ssetmatrix( n-i, nb, work, 0, ldwork, dwork, dwork_offset, lddwork, queue );
magma_ssyr2k(MagmaLower, MagmaNoTrans, n-i-nb, nb, c_neg_one,
dA(i+nb, i), ldda,
dwork, (dwork_offset+nb), lddwork, d_one,
dA(i+nb, i+nb), ldda, queue);
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
MAGMA_S_SET2REAL( *A(j+1, j), e[j] );
d[j] = MAGMA_S_REAL( *A(j, j) );
}
}
#ifdef FAST_SYMV
magma_free( dwork2 );
#endif
/* Use unblocked code to reduce the last or only block */
if (1<=n-nx)
magma_sgetmatrix( n-i, n-i, dA(i, i), ldda, A(i, i), 0, lda, queue );
i_n = n-i;
lapackf77_ssytrd(uplo_, &i_n, A(i, i), &lda, &d[i], &e[i],
&tau[i], work, &lwork, &iinfo);
}
magma_free( da );
// ACD
// MAGMA_S_SET2REAL( work[0], lwkopt );
MAGMA_S_SET2REAL( work[0], (float) lwkopt );
return *info;
} /* magma_ssytrd */
/**
Purpose
-------
SGEQRF computes a QR factorization of a REAL M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
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,out]
A REAL array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau REAL array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= max( N*NB, 2*NB*NB ),
where NB can be obtained through magma_get_sgeqrf_nb(M).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
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-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sgeqrf(magma_int_t m, magma_int_t n,
float *A, magma_int_t lda, float *tau,
float *work, magma_int_t lwork,
magma_int_t *info )
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
float *dA, *dwork, *dT;
float c_one = MAGMA_S_ONE;
magma_int_t i, k, lddwork, old_i, old_ib;
magma_int_t ib, ldda;
/* Function Body */
*info = 0;
magma_int_t nb = magma_get_sgeqrf_nb(min(m, n));
// need 2*nb*nb to store T and upper triangle of V simultaneously
magma_int_t lwkopt = max(n*nb, 2*nb*nb);
work[0] = MAGMA_S_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1, lwkopt) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
// largest N for larfb is n-nb (trailing matrix lacks 1st panel)
lddwork = ((n+31)/32)*32 - nb;
ldda = ((m+31)/32)*32;
magma_int_t num_gpus = magma_num_gpus();
if ( num_gpus > 1 ) {
/* call multiple-GPU interface */
return magma_sgeqrf4(num_gpus, m, n, A, lda, tau, work, lwork, info);
}
// allocate space for dA, dwork, and dT
if (MAGMA_SUCCESS != magma_smalloc( &dA, n*ldda + nb*lddwork + nb*nb )) {
/* Switch to the "out-of-core" (out of GPU-memory) version */
return magma_sgeqrf_ooc(m, n, A, lda, tau, work, lwork, info);
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = current_stream;
}
dwork = dA + n*ldda;
dT = dA + n*ldda + nb*lddwork;
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
Asynchronously send the matrix to the GPU except the first panel. */
magma_ssetmatrix_async( m, n-nb,
A(0,nb), lda,
dA(0,nb), ldda, stream[0] );
old_i = 0;
old_ib = nb;
for (i = 0; i < k-nb; i += nb) {
ib = min(k-i, nb);
if (i > 0) {
/* download i-th panel */
magma_queue_sync( stream[1] );
magma_sgetmatrix_async( m-i, ib,
dA(i,i), ldda,
A(i,i), lda, stream[0] );
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i), ldda, dT, nb,
dA(old_i, old_i+2*old_ib), ldda, dwork, lddwork);
magma_sgetmatrix_async( i, ib,
dA(0,i), ldda,
A(0,i), lda, stream[1] );
magma_queue_sync( stream[0] );
}
magma_int_t rows = m-i;
lapackf77_sgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(i,i), &lda, tau+i, work, &ib);
spanel_to_q(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
/* download the i-th V matrix */
magma_ssetmatrix_async( rows, ib, A(i,i), lda, dA(i,i), ldda, stream[0] );
/* download the T matrix */
magma_queue_sync( stream[1] );
magma_ssetmatrix_async( ib, ib, work, ib, dT, nb, stream[0] );
magma_queue_sync( stream[0] );
if (i + ib < n) {
if (i+ib < k-nb) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left (look-ahead) */
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dT, nb,
dA(i, i+ib), ldda, dwork, lddwork);
sq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
else {
/* After last panel, update whole trailing matrix. */
/* Apply H' to A(i:m,i+ib:n) from the left */
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dT, nb,
dA(i, i+ib), ldda, dwork, lddwork);
sq_to_panel(MagmaUpper, ib, A(i,i), lda, work+ib*ib);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
if (i != 0) {
magma_sgetmatrix_async( m, ib, dA(0,i), ldda, A(0,i), lda, stream[1] );
magma_queue_sync( stream[1] );
}
magma_int_t rows = m-i;
lapackf77_sgeqrf(&rows, &ib, A(i,i), &lda, tau+i, work, &lwork, info);
}
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
magma_free( dA );
return *info;
} /* magma_sgeqrf */
/**
Purpose
-------
DGEQLF computes a QL factorization of a DOUBLE_PRECISION M-by-N matrix A:
A = Q * L.
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,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, if m >= n, the lower triangle of the subarray
A(m-n+1:m,1:n) contains the N-by-N lower triangular matrix L;
if m <= n, the elements on and below the (n-m)-th
superdiagonal contain the M-by-N lower trapezoidal matrix L;
the remaining elements, with the array TAU, represent the
orthogonal matrix Q as a product of elementary reflectors
(see Further Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau DOUBLE_PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= max(1,N,2*NB^2).
For optimum performance LWORK >= max(N*NB, 2*NB^2) where NB can be obtained
through magma_get_dgeqlf_nb(M).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(k) . . . H(2) H(1), where k = min(m,n).
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(m-k+i+1:m) = 0 and v(m-k+i) = 1; v(1:m-k+i-1) is stored on exit in
A(1:m-k+i-1,n-k+i), and tau in TAU(i).
@ingroup magma_dgeqlf_comp
********************************************************************/
extern "C" magma_int_t
magma_dgeqlf(
magma_int_t m, magma_int_t n,
double *A, magma_int_t lda, double *tau,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dwork(i_) (dwork + (i_))
magmaDouble_ptr dA, dwork;
double c_one = MAGMA_D_ONE;
magma_int_t i, k, lddwork, old_i, old_ib, nb;
magma_int_t rows, cols;
magma_int_t ib, ki, kk, mu, nu, iinfo, ldda;
int lquery;
nb = magma_get_dgeqlf_nb(m);
*info = 0;
lquery = (lwork == -1);
// silence "uninitialized" warnings
old_ib = nb;
old_i = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
}
k = min(m,n);
if (*info == 0) {
if (k == 0)
work[0] = c_one;
else {
work[0] = MAGMA_D_MAKE( max(n*nb, 2*nb*nb), 0 );
}
if (lwork < max(max(1,n), 2*nb*nb) && ! lquery)
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (k == 0)
return *info;
lddwork = ((n+31)/32)*32;
ldda = ((m+31)/32)*32;
if (MAGMA_SUCCESS != magma_dmalloc( &dA, (n)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork = dA + ldda*n;
magma_queue_t queues[2];
magma_queue_create( &queues[0] );
magma_queue_create( &queues[1] );
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
The last kk columns are handled by the block method.
First, copy the matrix on the GPU except the last kk columns */
magma_dsetmatrix_async( m, n-nb,
A(0, 0), lda,
dA(0, 0), ldda, queues[0] );
ki = ((k - nb - 1) / nb) * nb;
kk = min(k, ki + nb);
for (i = k - kk + ki; i >= k -kk; i -= nb) {
ib = min(k-i,nb);
if (i < k - kk + ki) {
/* 1. Copy asynchronously the current panel to the CPU.
2. Copy asynchronously the submatrix below the panel
to the CPU) */
rows = m - k + i + ib;
magma_dgetmatrix_async( rows, ib,
dA(0, n-k+i), ldda,
A(0, n-k+i), lda, queues[1] );
magma_dgetmatrix_async( m-rows, ib,
dA(rows, n-k+i), ldda,
A(rows, n-k+i), lda, queues[0] );
/* Apply H' to A(1:m-k+i+ib-1,1:n-k+i-1) from the left in
two steps - implementing the lookahead techniques.
This is the main update from the lookahead techniques. */
rows = m - k + old_i + old_ib;
cols = n - k + old_i - old_ib;
magma_dlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, cols, old_ib,
dA(0, cols+old_ib), ldda, dwork(0), lddwork,
dA(0, 0 ), ldda, dwork(old_ib), lddwork);
}
magma_queue_sync( queues[1] );
/* Compute the QL factorization of the current block
A(1:m-k+i+ib-1,n-k+i:n-k+i+ib-1) */
rows = m - k + i + ib;
cols = n - k + i;
lapackf77_dgeqlf( &rows, &ib, A(0,cols), &lda, tau+i, work, &lwork, &iinfo );
if (cols > 0) {
/* Form the triangular factor of the block reflector
H = H(i+ib-1) . . . H(i+1) H(i) */
lapackf77_dlarft( MagmaBackwardStr, MagmaColumnwiseStr,
&rows, &ib,
A(0, cols), &lda, tau + i, work, &ib);
dpanel_to_q( MagmaLower, ib, A(rows-ib,cols), lda, work+ib*ib);
magma_dsetmatrix( rows, ib,
A(0,cols), lda,
dA(0,cols), ldda );
dq_to_panel( MagmaLower, ib, A(rows-ib,cols), lda, work+ib*ib);
// Send the triangular part on the GPU
magma_dsetmatrix( ib, ib, work, ib, dwork(0), lddwork );
/* Apply H' to A(1:m-k+i+ib-1,1:n-k+i-1) from the left in
two steps - implementing the lookahead techniques.
This is the update of first ib columns. */
if (i-ib >= k -kk)
magma_dlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, ib, ib,
dA(0, cols), ldda, dwork(0), lddwork,
dA(0,cols-ib), ldda, dwork(ib), lddwork);
else {
magma_dlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, cols, ib,
dA(0, cols), ldda, dwork(0), lddwork,
dA(0, 0 ), ldda, dwork(ib), lddwork);
}
old_i = i;
old_ib = ib;
}
}
mu = m - k + i + nb;
nu = n - k + i + nb;
magma_dgetmatrix( m, nu, dA(0,0), ldda, A(0,0), lda );
} else {
mu = m;
nu = n;
}
/* Use unblocked code to factor the last or only block */
if (mu > 0 && nu > 0)
lapackf77_dgeqlf(&mu, &nu, A(0,0), &lda, tau, work, &lwork, &iinfo);
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_dgeqlf */
extern "C" magma_int_t
magma_dgeqp3( magma_int_t m, magma_int_t n,
double *A, magma_int_t lda,
magma_int_t *jpvt, double *tau,
double *work, magma_int_t lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
double *rwork,
#endif
magma_int_t *info )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DGEQP3 computes a QR factorization with column pivoting of a
matrix A: A*P = Q*R using Level 3 BLAS.
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.
A (input/output) DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the upper triangle of the array contains the
min(M,N)-by-N upper trapezoidal matrix R; the elements below
the diagonal, together with the array TAU, represent the
unitary matrix Q as a product of min(M,N) elementary
reflectors.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
JPVT (input/output) INTEGER array, dimension (N)
On entry, if JPVT(J).ne.0, the J-th column of A is permuted
to the front of A*P (a leading column); if JPVT(J)=0,
the J-th column of A is a free column.
On exit, if JPVT(J)=K, then the J-th column of A*P was the
the K-th column of A.
TAU (output) DOUBLE_PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors.
WORK (workspace/output) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO=0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK.
For [sd]geqp3, LWORK >= (N+1)*NB + 2*N;
for [cz]geqp3, LWORK >= (N+1)*NB,
where NB is the optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
For [cz]geqp3 only:
RWORK (workspace) DOUBLE PRECISION array, dimension (2*N)
INFO (output) INTEGER
= 0: successful exit.
< 0: if INFO = -i, the i-th argument had an illegal value.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
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-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in
A(i+1:m,i), and tau in TAU(i).
===================================================================== */
#define A(i, j) (A + (i) + (j)*(lda ))
#define dA(i, j) (dwork + (i) + (j)*(ldda))
double *dwork, *df;
magma_int_t ione = 1;
magma_int_t n_j, ldda, ldwork;
magma_int_t j, jb, na, nb, sm, sn, fjb, nfxd, minmn;
magma_int_t topbmn, sminmn, lwkopt, lquery;
*info = 0;
lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
}
nb = magma_get_dgeqp3_nb(min(m, n));
if (*info == 0) {
minmn = min(m,n);
if (minmn == 0) {
lwkopt = 1;
} else {
lwkopt = (n + 1)*nb;
#if defined(PRECISION_d) || defined(PRECISION_s)
lwkopt += 2*n;
#endif
}
work[0] = MAGMA_D_MAKE( lwkopt, 0. );
if (lwork < lwkopt && ! lquery) {
*info = -8;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
} else if (lquery) {
return *info;
}
if (minmn == 0)
return *info;
#if defined(PRECISION_d) || defined(PRECISION_s)
double *rwork = work + (n + 1)*nb;
#endif
ldda = ((m+31)/32)*32;
ldwork = n*ldda + (n+1)*nb;
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
df = dwork + n*ldda;
// dwork used for dA
magma_queue_t stream;
magma_queue_create( &stream );
/* Move initial columns up front.
* Note jpvt uses 1-based indices for historical compatibility. */
nfxd = 0;
for (j = 0; j < n; ++j) {
if (jpvt[j] != 0) {
if (j != nfxd) {
blasf77_dswap(&m, A(0, j), &ione, A(0, nfxd), &ione);
jpvt[j] = jpvt[nfxd];
jpvt[nfxd] = j + 1;
}
else {
jpvt[j] = j + 1;
}
++nfxd;
}
else {
jpvt[j] = j + 1;
}
}
/* Factorize fixed columns
=======================
Compute the QR factorization of fixed columns and update
remaining columns. */
if (nfxd > 0) {
na = min(m,nfxd);
lapackf77_dgeqrf(&m, &na, A, &lda, tau, work, &lwork, info);
if (na < n) {
n_j = n - na;
lapackf77_dormqr( MagmaLeftStr, MagmaTransStr, &m, &n_j, &na,
A, &lda, tau, A(0, na), &lda,
work, &lwork, info );
}
}
/* Factorize free columns */
if (nfxd < minmn) {
sm = m - nfxd;
sn = n - nfxd;
sminmn = minmn - nfxd;
if (nb < sminmn) {
j = nfxd;
// Set the original matrix to the GPU
magma_dsetmatrix_async( m, sn,
A (0,j), lda,
dA(0,j), ldda, stream );
}
/* Initialize partial column norms. */
for (j = nfxd; j < n; ++j) {
rwork[j] = cblas_dnrm2(sm, A(nfxd, j), ione);
rwork[n + j] = rwork[j];
}
j = nfxd;
if (nb < sminmn) {
/* Use blocked code initially. */
magma_queue_sync( stream );
/* Compute factorization: while loop. */
topbmn = minmn - nb;
while(j < topbmn) {
jb = min(nb, topbmn - j);
/* Factorize JB columns among columns J:N. */
n_j = n - j;
if (j>nfxd) {
// Get panel to the CPU
magma_dgetmatrix( m-j, jb,
dA(j,j), ldda,
A (j,j), lda );
// Get the rows
magma_dgetmatrix( jb, n_j - jb,
dA(j,j + jb), ldda,
A (j,j + jb), lda );
}
magma_dlaqps( m, n_j, j, jb, &fjb,
A (0, j), lda,
dA(0, j), ldda,
&jpvt[j], &tau[j], &rwork[j], &rwork[n + j],
work,
&work[jb], n_j,
&df[jb], n_j );
j += fjb; /* fjb is actual number of columns factored */
}
}
/* Use unblocked code to factor the last or only block. */
if (j < minmn) {
n_j = n - j;
if (j > nfxd) {
magma_dgetmatrix( m-j, n_j,
dA(j,j), ldda,
A (j,j), lda );
}
lapackf77_dlaqp2(&m, &n_j, &j, A(0, j), &lda, &jpvt[j],
&tau[j], &rwork[j], &rwork[n+j], work );
}
}
work[0] = MAGMA_D_MAKE( lwkopt, 0. );
magma_free( dwork );
magma_queue_destroy( stream );
return *info;
} /* dgeqp3 */
/**
Purpose
-------
CLAHRU is an auxiliary MAGMA routine that is used in CGEHRD to update
the trailing sub-matrices after the reductions of the corresponding
panels.
See further details below.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ihi INTEGER
Last row to update. Same as IHI in cgehrd.
@param[in]
k INTEGER
Number of rows of the matrix Am (see details below)
@param[in]
nb INTEGER
Block size
@param[out]
A COMPLEX array, dimension (LDA,N-K)
On entry, the N-by-(N-K) general matrix to be updated. The
computation is done on the GPU. After Am is updated on the GPU
only Am(1:NB) is transferred to the CPU - to update the
corresponding Am matrix. See Further Details below.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
This implementation follows the 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.
The difference is that here Am is computed on the GPU.
M is renamed Am, G is renamed Ag.
@ingroup magma_cgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_clahru_m(
magma_int_t n, magma_int_t ihi, magma_int_t k, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
struct cgehrd_data* data )
{
#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 dW( d, i, j ) (data->W [d] + (i) + (j)*ldda)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
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 d;
magma_int_t dk, dkhi, dknb, dn;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (ihi < 0 || ihi > n) {
info = -2;
} else if (k < 0 || k > n) {
info = -3;
} else if (nb < 1 || nb > n) {
info = -4;
} else if (lda < max(1,n)) {
info = -6;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k, ihi, &dk, &dkhi );
magma_indices_1D_bcyclic( nb, ngpu, d, k+nb, n, &dknb, &dn );
// -----
// on right, A := A Q = A - A V T V'
// Update Am = Am - Am V T Vd' = Am - Ym Wd', with Wd = Vd T'
// Wd = Vd T' = V(k:ihi-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
// Vd and Wd are the portions corresponding to the block cyclic dkstribution
magma_cgemm( MagmaNoTrans, MagmaConjTrans, dkhi-dk, nb, nb,
c_one, dVd(d, dk, 0), ldvd,
dTi(d), nb,
c_zero, dW (d, dk, 0), ldda );
// Am = Am - Ym Wd' = A(0:k-1, k:ihi-1) - Ym(0:k-1, 0:nb-1) * W(k:ihi-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, k, dkhi-dk, nb,
c_neg_one, dY(d, 0, 0), ldda,
dW(d, dk, 0), ldda,
c_one, dA(d, 0, dk), ldda );
// -----
// on right, A := A Q = A - A V T V'
// Update Ag = Ag - Ag V T V' = Ag - Yg Wd'
// Ag = Ag - Yg Wd' = A(k:ihi-1, nb:ihi-k-1) - Y(k:ihi-1, 0:nb-1) * W(k+nb:ihi-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, dkhi-dknb, nb,
c_neg_one, dY(d, k, 0), ldda,
dW(d, dknb, 0), ldda,
c_one, dA(d, k, dknb), ldda );
// -----
// on left, A := Q' A = A - V T' V' A
// Ag2 = Ag2 - V T' V' Ag2 = W Yg, with W = V T' and Yg = V' Ag2
// Note that Ag is A(k:ihi, nb+1:ihi-k)
// while Ag2 is A(k:ihi, nb+1: n -k)
// here V and W are the whole matrices, not just block cyclic portion
// W = V T' = V(k:ihi-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
// TODO would it be cheaper to compute the whole matrix and
// copy the block cyclic portions to another workspace?
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, nb, nb,
c_one, dV (d, k, 0), ldv,
dTi(d), nb,
c_zero, dW (d, k, 0), ldda );
// Z = V(k:ihi-1, 0:nb-1)' * A(k:ihi-1, nb:n-k-1); Z is stored over Y
magma_cgemm( MagmaConjTrans, MagmaNoTrans, nb, dn-dknb, ihi-k,
c_one, dV(d, k, 0), ldv,
dA(d, k, dknb), ldda,
c_zero, dY(d, 0, 0), nb );
// Ag2 = Ag2 - W Z = A(k:ihi-1, k+nb:n-1) - W(k+nb:n-1, 0:nb-1) * Z(0:nb-1, k+nb:n-1)
magma_cgemm( MagmaNoTrans, MagmaNoTrans, ihi-k, dn-dknb, nb,
c_neg_one, dW(d, k, 0), ldda,
dY(d, 0, 0), nb,
c_one, dA(d, k, dknb), ldda );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return info;
}
/**
Purpose
-------
SSYTRD reduces a real symmetric matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
Arguments
---------
@param[in]
num_gpus INTEGER
The number of GPUs. num_gpus > 0.
@param[in]
num_streams INTEGER
The number of GPU streams used for update. 10 >= num_streams > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = MagmaUpper, the diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, and the elements below the first subdiagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
d REAL array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
tau REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_ssyev_comp
********************************************************************/
extern "C" magma_int_t
magma_ssytrd_mgpu(
magma_int_t num_gpus, magma_int_t num_streams, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *d, float *e, float *tau,
float *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define dA(id, i, j) (dA[(id)] + (j)*ldda + (i))
#define dW(id, i, j) (dwork[(id)] + (j)*ldda + (i))
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ln, ldda;
magma_int_t nb = magma_get_ssytrd_nb(n), ib;
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float d_one = MAGMA_D_ONE;
//float mv_time = 0.0;
#ifdef PROFILE_SY2RK
float up_time = 0.0;
#endif
magma_int_t kk, nx;
magma_int_t i = 0, ii, iii, j, did, i_n;
magma_int_t iinfo;
magma_int_t ldwork, lddwork, lwkopt, ldwork2;
magma_int_t lquery;
magma_queue_t stream[MagmaMaxGPUs][10];
float *dx[MagmaMaxGPUs], *dy[MagmaMaxGPUs], *hwork;
float *dwork2[MagmaMaxGPUs];
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
} else if ( num_streams > 2 ) {
*info = 2; // TODO fix
}
/* Determine the block size. */
ldwork = lddwork = n;
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_S_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
float *dA[MagmaMaxGPUs];
float *dwork[MagmaMaxGPUs];
float times[11];
for( did=0; did < 11; did++ )
times[did] = 0;
//#define PROFILE_SY2RK
#ifdef PROFILE_SY2RK
magma_event_t start, stop;
float etime;
magma_setdevice(0);
magma_event_create( &start );
magma_event_create( &stop );
#endif
ldda = lda;
ln = ((nb*(1+n/(nb*num_gpus))+31)/32)*32;
ldwork2 = (1+ n / nb + (n % nb != 0)) * ldda;
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
// TODO fix memory leak
if ( MAGMA_SUCCESS != magma_smalloc(&dA[did], ln*ldda+3*lddwork*nb) ||
MAGMA_SUCCESS != magma_smalloc(&dx[did], num_streams*n) ||
MAGMA_SUCCESS != magma_smalloc(&dy[did], num_streams*n) ||
MAGMA_SUCCESS != magma_smalloc(&dwork2[did], ldwork2 ) ) {
for( i=0; i < did; i++ ) {
magma_setdevice(i);
magma_free(dA[i]);
magma_free(dx[i]);
magma_free(dy[i]);
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork[did] = dA[did] + ln*ldda;
for( kk=0; kk < num_streams; kk++ )
magma_queue_create(&stream[did][kk]);
}
magma_setdevice(0);
// TODO fix memory leak dwork2
if ( MAGMA_SUCCESS != magma_smalloc_pinned( &hwork, num_streams*num_gpus*n ) ) {
for( i=0; i < num_gpus; i++ ) {
magma_setdevice(i);
magma_free(dA[i]);
magma_free(dx[i]);
magma_free(dy[i]);
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (n < 2048)
nx = n;
else
nx = 512;
if (upper) {
/* Copy the matrix to the GPU */
if (1 <= n-nx) {
magma_shtodhe(num_gpus, uplo, n, nb, A, lda, dA, ldda, stream, &iinfo );
}
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
for (i = nb*((n-1)/nb); i >= nx; i -= nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
/* wait for the next panel */
if (i != nb*((n-1)/nb)) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
magma_slatrd_mgpu(num_gpus, uplo, n, i+ib, ib, nb,
A(0, 0), lda, e, tau,
work, ldwork,
dA, ldda, 0,
dwork, i+ib,
dwork2, ldwork2,
1, dx, dy, hwork,
stream, times);
magma_ssyr2k_mgpu(num_gpus, MagmaUpper, MagmaNoTrans, nb, i, ib,
c_neg_one, dwork, i+ib, 0,
d_one, dA, ldda, 0,
num_streams, stream);
/* get the next panel */
if (i-nb >= nx ) {
ib = min(nb, n-(i-nb));
ii = nb*((i-nb)/(nb*num_gpus));
did = ((i-nb)/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( (i-nb)+ib, ib,
dA(did, 0, ii), ldda,
A(0, i-nb), lda,
stream[did][0] );
}
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+ib; ++j) {
if ( j > 0 ) {
*A(j-1,j) = MAGMA_S_MAKE( e[j - 1], 0 );
}
d[j] = MAGMA_S_REAL( *A(j, j) );
}
} /* end of for i=... */
if ( nx > 0 ) {
if (1 <= n-nx) { /* else A is already on CPU */
for (i=0; i < nx; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( nx, ib,
dA(did, 0, ii), ldda,
A(0, i), lda,
stream[did][0] );
}
}
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
/* Use unblocked code to reduce the last or only block */
lapackf77_ssytd2(uplo_, &nx, A(0, 0), &lda, d, e, tau, &iinfo);
}
}
else {
trace_init( 1, num_gpus, num_streams, (CUstream_st**)stream );
/* Copy the matrix to the GPU */
if (1 <= n-nx) {
magma_shtodhe(num_gpus, uplo, n, nb, A, lda, dA, ldda, stream, &iinfo );
}
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
/* Reduce columns i:i+ib-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != 0) {
magma_setdevice(did);
trace_gpu_start( did, 0, "comm", "get" );
magma_sgetmatrix_async( n-i, ib,
dA(did, i, ii), ldda,
A(i,i), lda,
stream[did][0] );
trace_gpu_end( did, 0 );
magma_queue_sync(stream[did][0]);
magma_setdevice(0);
}
magma_slatrd_mgpu(num_gpus, uplo, n, n-i, ib, nb,
A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA, ldda, i,
dwork, (n-i),
dwork2, ldwork2,
1, dx, dy, hwork,
stream, times );
#ifdef PROFILE_SY2RK
magma_setdevice(0);
if ( i > 0 ) {
cudaEventElapsedTime(&etime, start, stop);
up_time += (etime/1000.0);
}
magma_event_record(start, 0);
#endif
magma_ssyr2k_mgpu(num_gpus, MagmaLower, MagmaNoTrans, nb, n-i-ib, ib,
c_neg_one, dwork, n-i, ib,
d_one, dA, ldda, i+ib, num_streams, stream);
#ifdef PROFILE_SY2RK
magma_setdevice(0);
magma_event_record(stop, 0);
#endif
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+ib; ++j) {
if ( j+1 < n ) {
*A(j+1,j) = MAGMA_S_MAKE( e[j], 0 );
}
d[j] = MAGMA_S_REAL( *A(j, j) );
}
} /* for i=... */
/* Use unblocked code to reduce the last or only block */
if ( i < n ) {
iii = i;
i_n = n-i;
if ( i > 0 ) {
for (; i < n; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( i_n, ib,
dA(did, iii, ii), ldda,
A(iii, i), lda,
stream[did][0] );
}
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
}
lapackf77_ssytrd(uplo_, &i_n, A(iii, iii), &lda, &d[iii], &e[iii],
&tau[iii], work, &lwork, &iinfo);
}
}
#ifdef PROFILE_SY2RK
magma_setdevice(0);
if ( n > nx ) {
cudaEventElapsedTime(&etime, start, stop);
up_time += (etime/1000.0);
}
magma_event_destroy( start );
magma_event_destroy( stop );
#endif
trace_finalize( "ssytrd.svg", "trace.css" );
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
for( kk=0; kk < num_streams; kk++ )
magma_queue_sync(stream[did][kk]);
for( kk=0; kk < num_streams; kk++ )
magma_queue_destroy(stream[did][kk]);
magma_free(dA[did]);
magma_free(dx[did]);
magma_free(dy[did]);
magma_free(dwork2[did]);
}
magma_free_pinned(hwork);
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
work[0] = MAGMA_S_MAKE( lwkopt, 0 );
#ifdef PROFILE_SY2RK
printf( " n=%d nb=%d\n", n, nb );
printf( " Time in SLARFG: %.2e seconds\n", times[0] );
//printf( " Time in SSYMV : %.2e seconds\n", mv_time );
printf( " Time in SSYR2K: %.2e seconds\n", up_time );
#endif
return *info;
} /* magma_ssytrd */
