extern "C" magma_int_t
magma_svinit(
magma_s_matrix *x,
magma_location_t mem_loc,
magma_int_t num_rows,
magma_int_t num_cols,
float values,
magma_queue_t queue )
{
magma_int_t info = 0;
x->val = NULL;
x->diag = NULL;
x->row = NULL;
x->rowidx = NULL;
x->col = NULL;
x->list = NULL;
x->blockinfo = NULL;
x->dval = NULL;
x->ddiag = NULL;
x->drow = NULL;
x->drowidx = NULL;
x->dcol = NULL;
x->dlist = NULL;
x->storage_type = Magma_DENSE;
x->memory_location = mem_loc;
x->sym = Magma_GENERAL;
x->diagorder_type = Magma_VALUE;
x->fill_mode = MagmaFull;
x->num_rows = num_rows;
x->num_cols = num_cols;
x->nnz = num_rows*num_cols;
x->max_nnz_row = num_cols;
x->diameter = 0;
x->blocksize = 1;
x->numblocks = 1;
x->alignment = 1;
x->major = MagmaColMajor;
x->ld = num_rows;
if ( mem_loc == Magma_CPU ) {
CHECK( magma_smalloc_cpu( &x->val, x->nnz ));
for( magma_int_t i=0; i<x->nnz; i++) {
x->val[i] = values;
}
}
else if ( mem_loc == Magma_DEV ) {
CHECK( magma_smalloc( &x->val, x->nnz ));
magmablas_slaset( MagmaFull, x->num_rows, x->num_cols, values, values, x->val, x->num_rows, queue );
}
cleanup:
return info;
}
/**
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.
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 REAL 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 SGEQRF_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 REAL array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by SGEQRF_GPU.
@param[in]
T REAL array, dimension (NB, min(M,N)).
T contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_sgeqrf_gpu (except stored on the CPU, not the GPU).
@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 T.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sorgqr_m(
magma_int_t m, magma_int_t n, magma_int_t k,
float *A, magma_int_t lda,
float *tau,
float *T, magma_int_t nb,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(d,i,j) (dA[d] + (i) + (j)*ldda)
#define dT(d,i,j) (dT[d] + (i) + (j)*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldwork;
magma_int_t d, i, ib, j, jb, ki, kk;
float *work=NULL;
*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;
}
magma_int_t di, dn;
magma_int_t dpanel;
magma_int_t ngpu = magma_num_gpus();
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// Allocate memory on GPUs for A and workspaces
magma_int_t ldda = magma_roundup( m, 32 );
magma_int_t lddwork = magma_roundup( n, 32 );
magma_int_t min_lblocks = (n / nb) / ngpu; // min. blocks per gpu
magma_int_t last_dev = (n / nb) % ngpu; // device with last block
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
float *dA[ MagmaMaxGPUs ] = { NULL };
float *dT[ MagmaMaxGPUs ] = { NULL };
float *dV[ MagmaMaxGPUs ] = { NULL };
float *dW[ MagmaMaxGPUs ] = { NULL };
magma_queue_t queues[ MagmaMaxGPUs ] = { NULL };
for( d = 0; d < ngpu; ++d ) {
// example with n = 75, nb = 10, ngpu = 3
// min_lblocks = 2
// last_dev = 1
// gpu 0: 2 blocks, cols: 0- 9, 30-39, 60-69
// gpu 1: 1+ blocks, cols: 10-19, 40-49, 70-74 (partial)
// gpu 2: 1 block, cols: 20-29, 50-59
magma_setdevice( d );
nlocal[d] = min_lblocks*nb;
if ( d < last_dev ) {
nlocal[d] += nb;
}
else if ( d == last_dev ) {
nlocal[d] += (n % nb);
}
ldwork = nlocal[d]*ldda // dA
+ nb*m // dT
+ nb*ldda // dV
+ nb*lddwork; // dW
if ( MAGMA_SUCCESS != magma_smalloc( &dA[d], ldwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
goto cleanup;
}
dT[d] = dA[d] + nlocal[d]*ldda;
dV[d] = dT[d] + nb*m;
dW[d] = dV[d] + nb*ldda;
magma_queue_create( d, &queues[d] );
}
trace_init( 1, ngpu, 1, queues );
// 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 CPU work space
// n*nb for larfb work
// m*nb for V
// nb*nb for T
lwork = (n + m + nb) * nb;
magma_smalloc_cpu( &work, lwork );
if (work == NULL) {
*info = MAGMA_ERR_HOST_ALLOC;
goto cleanup;
}
float *work_T, *work_V;
work_T = work + n*nb;
work_V = work + n*nb + nb*nb;
// Use unblocked code for the last or only block.
if (kk < n) {
trace_cpu_start( 0, "ungqr", "ungqr last block" );
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
// sorgqr requires less workspace (n*nb), but is slow if k < sorgqr's block size.
// replacing it with the 4 routines below is much faster (e.g., 60x).
//magma_int_t iinfo;
//lapackf77_sorgqr( &m_kk, &n_kk, &k_kk,
// A(kk, kk), &lda,
// &tau[kk], work, &lwork, &iinfo );
lapackf77_slacpy( MagmaFullStr, &m_kk, &k_kk, A(kk,kk), &lda, work_V, &m_kk);
lapackf77_slaset( MagmaFullStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
work_V, &m_kk, &tau[kk], work_T, &k_kk);
lapackf77_slarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
work_V, &m_kk, work_T, &k_kk, A(kk, kk), &lda, work, &n_kk );
if (kk > 0) {
for( j=kk; j < n; j += nb ) {
jb = min( n-j, nb );
d = (j / nb) % ngpu;
di = ((j / nb) / ngpu) * nb;
magma_setdevice( d );
magma_ssetmatrix( m_kk, jb,
A(kk, j), lda,
dA(d, kk, di), ldda, queues[d] );
// Set A(1:kk,kk+1:n) to zero.
magmablas_slaset( MagmaFull, kk, jb, c_zero, c_zero, dA(d, 0, di), ldda, queues[d] );
}
}
trace_cpu_end( 0 );
}
if (kk > 0) {
// Use blocked code
// send T to all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set T" );
magma_ssetmatrix_async( nb, min(m,n), T, nb, dT[d], nb, queues[d] );
trace_gpu_end( d, 0 );
}
// queue: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
for( i = ki; i >= 0; i -= nb ) {
ib = min(nb, k - i);
mi = m - i;
dpanel = (i / nb) % ngpu;
di = ((i / nb) / ngpu) * nb;
// Send current panel to dV on the GPUs
lapackf77_slaset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
trace_gpu_start( d, 0, "set", "set V" );
magma_ssetmatrix_async( mi, ib,
A(i, i), lda,
dV[d], ldda, queues[d] );
trace_gpu_end( d, 0 );
}
// set panel to identity
magma_setdevice( dpanel );
trace_gpu_start( dpanel, 0, "laset", "laset" );
magmablas_slaset( MagmaFull, i, ib, c_zero, c_zero, dA(dpanel, 0, di), ldda, queues[dpanel] );
magmablas_slaset( MagmaFull, mi, ib, c_zero, c_one, dA(dpanel, i, di), ldda, queues[dpanel] );
trace_gpu_end( dpanel, 0 );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_indices_1D_bcyclic( nb, ngpu, d, i, n, &di, &dn );
trace_gpu_start( d, 0, "larfb", "larfb" );
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, dn-di, ib,
dV[d], ldda, dT(d,0,i), nb,
dA(d, i, di), ldda, dW[d], lddwork, queues[d] );
trace_gpu_end( d, 0 );
}
}
}
// copy result back to CPU
trace_cpu_start( 0, "get", "get A" );
magma_sgetmatrix_1D_col_bcyclic( m, n, dA, ldda, A, lda, ngpu, nb, queues );
trace_cpu_end( 0 );
}
#ifdef TRACING
char name[80];
snprintf( name, sizeof(name), "sorgqr-n%d-ngpu%d.svg", m, ngpu );
trace_finalize( name, "trace.css" );
#endif
cleanup:
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_free( dA[d] );
magma_queue_destroy( queues[d] );
}
magma_free_cpu( work );
magma_setdevice( orig_dev );
return *info;
} /* magma_sorgqr */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing slaset
Code is very similar to testing_slacpy.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_R;
magmaFloat_ptr d_A;
float offdiag, diag;
magma_int_t M, N, size, lda, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("%% uplo M N offdiag diag CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("%%===================================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
for( int ival = 0; ival < 4; ++ival ) {
// test combinations of zero & non-zero:
// ival offdiag diag
// 0 0 0
// 1 0 3.14
// 2 1.23 0
// 3 1.23 3.14
offdiag = MAGMA_S_MAKE( 1.2345, 6.7890 ) * (ival / 2);
diag = MAGMA_S_MAKE( 3.1415, 2.7183 ) * (ival % 2);
M = opts.msize[itest];
N = opts.nsize[itest];
//M += 2; // space for insets
//N += 2;
lda = M;
ldda = magma_roundup( M, opts.align );
size = lda*N;
if ( uplo[iuplo] == MagmaLower ) {
// save lower trapezoid (with diagonal)
if ( M > N ) {
gbytes = sizeof(float) * (1.*M*N - 0.5*N*(N-1)) / 1e9;
} else {
gbytes = sizeof(float) * 0.5*M*(M+1) / 1e9;
}
}
else if ( uplo[iuplo] == MagmaUpper ) {
// save upper trapezoid (with diagonal)
if ( N > M ) {
gbytes = sizeof(float) * (1.*M*N - 0.5*M*(M-1)) / 1e9;
} else {
gbytes = sizeof(float) * 0.5*N*(N+1) / 1e9;
}
}
else {
// save entire matrix
gbytes = sizeof(float) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, float, size );
TESTING_MALLOC_CPU( h_R, float, size );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_S_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix( M, N, h_A, lda, d_A, ldda, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
//magmablas_slaset( uplo[iuplo], M-2, N-2, offdiag, diag, d_A+1+ldda, ldda, opts.queue ); // inset by 1 row & col
magmablas_slaset( uplo[iuplo], M, N, offdiag, diag, d_A, ldda, opts.queue );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//magma_int_t M2 = M-2; // inset by 1 row & col
//magma_int_t N2 = N-2;
//lapackf77_slaset( lapack_uplo_const( uplo[iuplo] ), &M2, &N2, &offdiag, &diag, h_A+1+lda, &lda );
lapackf77_slaset( lapack_uplo_const( uplo[iuplo] ), &M, &N, &offdiag, &diag, h_A, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if ( opts.verbose ) {
printf( "A= " ); magma_sprint( M, N, h_A, lda );
printf( "dA=" ); magma_sprint_gpu( M, N, d_A, ldda );
}
/* =====================================================================
Check the result
=================================================================== */
magma_sgetmatrix( M, N, d_A, ldda, h_R, lda, opts.queue );
blasf77_saxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &M, &N, h_R, &lda, work);
bool okay = (error == 0);
status += ! okay;
printf("%5s %5d %5d %9.4f %6.4f %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
lapack_uplo_const( uplo[iuplo] ), (int) M, (int) N,
real(offdiag), real(diag),
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(okay ? "ok" : "failed") );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
SPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
Auxiliary subroutine for spotrf2_ooc. It is multiple gpu interface to compute
Cholesky of a "rectangular" matrix.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_sposv_comp
********************************************************************/
extern "C" magma_int_t
magma_spotrf3_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb,
magmaFloat_ptr d_lA[], magma_int_t ldda,
magmaFloat_ptr d_lP[], magma_int_t lddp,
float *A, magma_int_t lda, magma_int_t h,
magma_queue_t queues[][3], magma_event_t events[][5],
magma_int_t *info )
{
#define Alo(i, j) (A + ((j)+off_j)*lda + (nb*(((i)/nb)%h)+off_i))
#define Aup(i, j) (A + (nb*(((j)/nb)%h)+off_j)*lda + (i+off_i))
#define dlA(id, i, j) (d_lA[(id)] + (j)*ldda + (i))
#define dlP(id, i, j, k) (d_lP[(id)] + (k)*nb*lddp + (j)*lddp + (i))
#define dlPT(id, i, j, k) (d_lP[(id)] + (k)*nb*lddp + (j)*nb + (i))
magma_int_t j, jb, nb0, nb2, d, dd, id, j_local, j_local2, buf;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float d_one = 1.0;
float d_neg_one = -1.0;
int upper = (uplo == MagmaUpper);
float *dlpanel;
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
const magma_int_t stream1 = 0, stream2 = 1, stream3 = 2;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (!upper && ngpu*ldda < max(1,n)) {
*info = -4;
} else if (upper && ldda < max(1,m)) {
*info = -4;
}
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 );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
/* used by strsm_work */
float c_zero = MAGMA_S_ZERO;
int trsm_nb = 128;
int trsm_n = trsm_nb*((nb+trsm_nb-1)/trsm_nb);
float *d_dinvA[MagmaMaxGPUs];
float *d_x[MagmaMaxGPUs];
#define dinvA(d,j) &(d_dinvA[(d)][(j)*trsm_nb*trsm_n])
#define dx(d,j) &(d_x[(d)][(j)*nb*m])
/*
* Allocate device memory for the inversed diagonal blocks, size=N*BLOCK_SIZE
*/
// TODO free memory on failure.
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
if ( (MAGMA_SUCCESS != magma_smalloc( &d_dinvA[d], 2*trsm_nb*trsm_n )) ||
(MAGMA_SUCCESS != magma_smalloc( &d_x[d], 2*nb*(upper ? n : m) )) ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
magma_setdevice(0);
#endif
/* initialization */
for( d=0; d < ngpu; d++ ) {
/* local-n and local-ld */
if (upper) {
n_local[d] = (n/(nb*ngpu))*nb;
if (d < (n/nb)%ngpu)
n_local[d] += nb;
else if (d == (n/nb)%ngpu)
n_local[d] += n%nb;
} else {
n_local[d] = (m/(nb*ngpu))*nb;
if (d < (m/nb)%ngpu)
n_local[d] += nb;
else if (d == (m/nb)%ngpu)
n_local[d] += m%nb;
}
}
/* == initialize the trace */
trace_init( 1, ngpu, 3, (CUstream_st**)queues );
if (upper) {
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
for (j=0; j < m; j += nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%ngpu;
buf = (j/nb)%ngpu; // right now, we have ngpu buffers, so id and buf are the same..
/* Set the local index where the current panel is */
j_local = j/(nb*ngpu);
jb = min(nb, (m-j));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if ( j > 0 ) {
magmablasSetKernelStream( queues[id][stream1] );
trace_gpu_start( id, stream1, "syrk", "syrk" );
magma_ssyrk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dlA(id, 0, nb*j_local), ldda,
d_one, dlA(id, j, nb*j_local), ldda);
trace_gpu_end( id, stream1 );
}
/* send the diagonal to cpu on stream1 */
trace_gpu_start( id, stream1, "comm", "D to CPU" );
magma_sgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
queues[id][stream1] );
trace_gpu_end( id, stream1 );
/* update off-diagonal blocks in the panel */
if ( j > 0 ) {
d = (j/nb+1)%ngpu;
for( dd=0; dd < ngpu; dd++ ) {
j_local2 = j_local+1;
if ( d > id ) j_local2 --;
nb0 = nb*j_local2; // number of local columns in the panel, while jb is panel-size (number of rows)
if ( n_local[d] > nb0 ) {
magma_setdevice(d);
magmablasSetKernelStream( queues[d][stream2] );
if ( d == id ) {
dlpanel = dlA(d,0,nb*j_local);
ldpanel = ldda;
// the GPU owns the row from start, and no need of synch.
//magma_queue_wait_event( queues[d][stream2], events[d][0] ); // rows arrived at gpu
magma_queue_wait_event( queues[d][stream2], events[d][4] ); // wait for look-ahead trsm to finish
} else {
dlpanel = dlP(d,nb,0,buf);
ldpanel = lddp;
magma_queue_wait_event( queues[d][stream2], events[d][0] ); // rows arrived at gpu
}
trace_gpu_start( d, stream2, "gemm", "gemm" );
magma_sgemm(MagmaConjTrans, MagmaNoTrans,
jb, n_local[d]-nb0, j,
c_neg_one, dlpanel, ldpanel,
dlA(d, 0, nb0), ldda,
c_one, dlA(d, j, nb0), ldda);
trace_gpu_end( d, stream2 );
magma_event_record( events[d][2], queues[d][stream2] );
}
d = (d+1)%ngpu;
}
}
/* wait for panel and factorize it on cpu */
magma_setdevice(id);
magma_queue_sync( queues[id][stream1] );
trace_cpu_start( 0, "getrf", "getrf" );
lapackf77_spotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus on stream1 */
if ( (j+jb) < n) {
d = (j/nb+1)%ngpu;
for( dd=0; dd < ngpu; dd++ ) {
if ( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
magma_setdevice(d);
trace_gpu_start( d, stream1, "comm", "comm" );
magma_ssetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, ldpanel,
queues[d][stream1] );
trace_gpu_end( d, stream1 );
magma_event_record( events[d][1], queues[d][stream1] );
d = (d+1)%ngpu;
}
} else {
magma_setdevice(id);
trace_gpu_start( id, stream1, "comm", "comm" );
magma_ssetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
queues[id][stream1] );
trace_gpu_end( id, stream1 );
}
/* panel-factorize the off-diagonal */
if ( (j+jb) < n) {
d = (j/nb+1)%ngpu;
for( dd=0; dd < ngpu; dd++ ) {
/* next column */
j_local2 = j_local+1;
if ( d > id ) j_local2--;
if ( d == id ) {
dlpanel = dlA(d,j,nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
nb2 = n_local[d] - j_local2*nb;
magma_setdevice(d);
if ( j+jb < m && d == (j/nb+1)%ngpu ) {
/* owns the next column, look-ahead next block on stream1 */
nb0 = min(nb, nb2);
magmablasSetKernelStream( queues[d][stream1] );
magma_queue_wait_event( queues[d][stream1], events[d][2] ); // wait for gemm update
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
magmablas_slaset( MagmaFull, trsm_nb, trsm_n, c_zero, c_zero, dinvA(d,0), trsm_nb );
magmablas_slaset( MagmaFull, nb0, jb, c_zero, c_zero, dx(d,0), nb0 );
magmablas_strsm_work( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_strsm( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
magma_event_record( events[d][4], queues[d][stream1] );
trace_gpu_end( d, stream1 );
} else if ( nb2 > 0 ) {
/* update all the blocks on stream2 */
magma_queue_wait_event( queues[d][stream2], events[d][1] ); // wait for cholesky factor
trace_gpu_start( d, stream2, "trsm", "trsm" );
magmablasSetKernelStream( queues[d][stream2] );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
magmablas_slaset( MagmaFull, trsm_nb, trsm_n, c_zero, c_zero, dinvA(d,0), trsm_nb );
magmablas_slaset( MagmaFull, nb2, jb, c_zero, c_zero, dx(d,0), nb2 );
magmablas_strsm_work( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_strsm( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream2 );
}
d = (d+1)%ngpu;
} /* end of for */
/* ========================================================== */
if ( j+jb < m ) {
d = (j/nb+1)%ngpu;
/* next column */
j_local2 = j_local+1;
if ( d > id ) j_local2--;
nb0 = min(nb, n_local[d]-nb*j_local2 );
/* even on 1 gpu, off-diagonals are copied to cpu (synchronize at the end). *
* so we have the Cholesky factor, but only diagonal submatrix of the big panel, *
* on cpu at the end. */
int d2, buf2;
magma_setdevice(d);
/* lookahead done */
magma_queue_wait_event( queues[d][stream3], events[d][4] );
trace_gpu_start( d, stream3, "comm", "row to CPU" );
magma_sgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
queues[d][stream3] );
trace_gpu_end( d, stream3 );
magma_event_record( events[d][3], queues[d][stream3] );
/* needed on pluto */
//magma_queue_sync( queues[d][stream3] );
/* broadcast rows to gpus on stream2 */
buf2 = ((j+jb)/nb)%ngpu;
for( d2=0; d2 < ngpu; d2++ ) {
if ( d2 != d ) {
magma_setdevice(d2);
trace_gpu_start( d2, stream3, "comm", "row to GPUs" );
magma_queue_wait_event( queues[d2][stream3], events[d][3] ); // rows arrived at cpu on stream3
magma_ssetmatrix_async( j+jb, nb0,
Aup(0,j+jb), lda,
dlP(d2,nb,0,buf2), lddp,
queues[d2][stream3] );
trace_gpu_end( d2, stream3 );
magma_event_record( events[d2][0], queues[d2][stream3] );
}
}
/* =========================== */
/* update the remaining blocks */
nb2 = n_local[d]-(nb*j_local2 + nb0);
if ( nb2 > 0 ) {
if ( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
magma_setdevice(d);
magmablasSetKernelStream( queues[d][stream2] );
trace_gpu_start( d, stream2, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
int flag = 0;
if (flag == 0) {
magma_queue_wait_event( queues[d][stream2], events[d][4] ); // lookahead -> diagonal inversion
} else {
magmablas_slaset( MagmaFull, trsm_nb, trsm_n, c_zero, c_zero, dinvA(d,flag), trsm_nb );
magma_queue_wait_event( queues[d][stream2], events[d][1] ); // panel received
}
magmablas_slaset( MagmaFull, nb2, jb, c_zero, c_zero, dx(d,1), nb2 );
magmablas_strsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
flag, dinvA(d,flag), dx(d,1) );
#else
magma_queue_wait_event( queues[d][stream2], events[d][1] ); // wait for cholesky factor
magma_strsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda);
#endif
trace_gpu_end( d, stream2 );
}
}
} /* end of strsm */
} /* end of for j=1, .., n */
} else {
/* ---------------------------------------------- */
/* Lower-triangular case */
/* > Compute the Cholesky factorization A = L*L'. */
/* ---------------------------------------------- */
for (j=0; j < n; j += nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%ngpu;
buf = (j/nb)%ngpu;
/* Set the local index where the current panel is */
j_local = j/(nb*ngpu);
jb = min(nb, (n-j));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if ( j > 0 ) {
magmablasSetKernelStream( queues[id][stream1] );
magma_ssyrk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda);
}
/* send the diagonal to cpu on stream1 */
magma_sgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
queues[id][stream1] );
/* update off-diagonal blocks of the panel */
if ( j > 0 ) {
d = (j/nb+1)%ngpu;
for( dd=0; dd < ngpu; dd++ ) {
j_local2 = j_local+1;
if ( d > id ) j_local2 --;
nb0 = nb*j_local2;
if ( nb0 < n_local[d] ) {
magma_setdevice(d);
magmablasSetKernelStream( queues[d][stream2] );
if ( d == id ) {
dlpanel = dlA(d, nb*j_local, 0);
ldpanel = ldda;
magma_queue_wait_event( queues[d][stream2], events[d][4] ); // wait for look-ahead trsm to finish
} else {
dlpanel = dlPT(d,0,nb,buf);
ldpanel = nb;
magma_queue_wait_event( queues[d][stream2], events[d][0] ); // rows arrived at gpu
}
magma_sgemm( MagmaNoTrans, MagmaConjTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, ldpanel,
c_one, dlA(d, nb0, j), ldda);
magma_event_record( events[d][2], queues[d][stream2] );
}
d = (d+1)%ngpu;
}
}
/* wait for the panel and factorized it on cpu */
magma_setdevice(id);
magma_queue_sync( queues[id][stream1] );
lapackf77_spotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus on stream1 */
if ( (j+jb) < m) {
d = (j/nb+1)%ngpu;
for( dd=0; dd < ngpu; dd++ ) {
if ( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
magma_setdevice(d);
magma_ssetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, ldpanel,
queues[d][stream1] );
magma_event_record( events[d][1], queues[d][stream1] );
d = (d+1)%ngpu;
}
} else {
magma_setdevice(id);
magma_ssetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
queues[id][stream1] );
}
/* panel factorize the off-diagonal */
if ( (j+jb) < m) {
d = (j/nb+1)%ngpu;
for( dd=0; dd < ngpu; dd++ ) {
/* next column */
j_local2 = j_local+1;
if ( d > id ) j_local2--;
if ( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2);
magma_setdevice(d);
if ( j+nb < n && d == (j/nb+1)%ngpu ) { /* owns next column, look-ahead next block on stream1 */
if ( j > 0 ) magma_queue_wait_event( queues[d][stream1], events[d][2] ); // wait for gemm update
magmablasSetKernelStream( queues[d][stream1] );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
magmablas_slaset( MagmaFull, trsm_nb, trsm_n, c_zero, c_zero, dinvA(d,0), trsm_nb );
magmablas_slaset( MagmaFull, nb0, jb, c_zero, c_zero, dx(d,0), nb0 );
magmablas_strsm_work( MagmaRight, MagmaLower,
MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_strsm( MagmaRight, MagmaLower,
MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
magma_event_record( events[d][4], queues[d][stream1] );
} else if ( nb2 > 0 ) { /* other gpus updating all the blocks on stream2 */
/* update the entire column */
magma_queue_wait_event( queues[d][stream2], events[d][1] ); // wait for the cholesky factor
magmablasSetKernelStream( queues[d][stream2] );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
magmablas_slaset( MagmaFull, trsm_nb, trsm_n, c_zero, c_zero, dinvA(d,0), trsm_nb );
magmablas_slaset( MagmaFull, nb2, jb, c_zero, c_zero, dx(d,0), nb2 );
magmablas_strsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_strsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
}
d = (d+1)%ngpu;
} /* end for d */
/* ========================================================== */
if ( j+jb < n ) {
d = (j/nb+1)%ngpu;
/* next column */
j_local2 = j_local+1;
if ( d > id ) j_local2--;
nb0 = min(nb, n_local[d]-nb*j_local2 );
/* even on 1 gpu, we copy off-diagonal to cpu (but don't synchronize). */
/* so we have the Cholesky factor on cpu at the end. */
int d2, buf2;
//#define SPOTRF_DEVICE_TO_DEVICE
#ifdef SPOTRF_DEVICE_TO_DEVICE
// lookahead done
/* broadcast the rows to gpus */
buf2 = ((j+jb)/nb)%ngpu;
for( d2=0; d2 < ngpu; d2++ ) {
magma_setdevice(d2);
magma_queue_wait_event( queues[d2][stream3], events[d][4] );
if ( d2 != d ) {
magma_scopymatrix_async( nb0, j+jb,
dlPT(d2,0,nb,buf2), nb, // first nbxnb reserved for diagonal block
dlA(d, nb*j_local2, 0), ldda,
queues[d2][stream3] );
magma_event_record( events[d2][0], queues[d2][stream3] );
} else {
magma_sgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
queues[d][stream3] );
}
}
#else
// lookahead done
magma_setdevice(d);
magma_queue_wait_event( queues[d][stream3], events[d][4] );
magma_sgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
queues[d][stream3] );
magma_event_record( events[d][3], queues[d][stream3] );
/* syn on rows on CPU, seem to be needed on Pluto */
//magma_queue_sync( queues[d][stream3] );
/* broadcast the rows to gpus */
buf2 = ((j+jb)/nb)%ngpu;
for( d2=0; d2 < ngpu; d2++ ) {
if ( d2 != d ) {
magma_setdevice(d2);
magma_queue_wait_event( queues[d2][stream3], events[d][3] ); // getmatrix done
magma_ssetmatrix_async( nb0, j+jb,
Alo(j+jb,0), lda,
dlPT(d2,0,nb,buf2), nb, // first nbxnb reserved for diagonal block
queues[d2][stream3] );
magma_event_record( events[d2][0], queues[d2][stream3] );
}
}
#endif
/* =================================== */
/* updates remaining blocks on stream2 */
nb2 = n_local[d] - (j_local2*nb + nb0);
if ( nb2 > 0 ) {
if ( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d,0,0,buf);
ldpanel = nb;
}
magma_setdevice(d);
magmablasSetKernelStream( queues[d][stream2] );
/* update the remaining blocks in the column */
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
int flag = 0;
if (flag == 0) {
magma_queue_wait_event( queues[d][stream2], events[d][4] ); // lookahead -> diagonal inversion
} else {
magmablas_slaset( MagmaFull, trsm_nb, trsm_n, c_zero, c_zero, dinvA(d,flag), trsm_nb );
magma_queue_wait_event( queues[d][stream2], events[d][1] ); // panel received
}
magmablas_slaset( MagmaFull, nb2, jb, c_zero, c_zero, dx(d,1), nb2 );
magmablas_strsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
flag, dinvA(d,flag), dx(d,1) );
#else
magma_queue_wait_event( queues[d][stream2], events[d][1] ); // panel received
magma_strsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda);
#endif
}
}
}
}
} /* end of else not upper */
/* == finalize the trace == */
trace_finalize( "spotrf.svg", "trace.css" );
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
for( j=0; j < 3; j++ ) {
magma_queue_sync( queues[d][j] );
}
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(STRSM_WORK)
magma_free( d_dinvA[d] );
magma_free( d_x[d] );
#endif
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_spotrf_mgpu */
/***************************************************************************//**
Purpose
-------
SGEHRD2 reduces a REAL general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H .
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ilo INTEGER
@param[in]
ihi INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to SGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, 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]
tau REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
@param[out]
work (workspace) REAL array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The length of the array WORK. LWORK >= max(1,N).
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
---------------
The matrix Q is represented as a product of (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-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, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
@verbatim
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
@endverbatim
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_gehrd
*******************************************************************************/
extern "C" magma_int_t
magma_sgehrd2(
magma_int_t n, magma_int_t ilo, magma_int_t ihi,
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)
#ifdef HAVE_clBLAS
#define dA(i_,j_) dwork, ((i_) + (j_)*ldda + nb*ldda*2)
#define dT(i_,j_) dT, ((i_) + (j_)*nb + dT_offset)
#define dV(i_,j_) dwork, ((i_) + (j_)*ldda + nb*ldda)
#define dwork(i_) dwork, ((i_))
#else
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dT(i_,j_) (dT + (i_) + (j_)*nb)
#define dV(i_,j_) (dV + (i_) + (j_)*ldda)
#define dwork(i_) (dwork + (i_))
#endif
// Constants
const float c_one = MAGMA_S_ONE;
const float c_zero = MAGMA_S_ZERO;
// Local variables
magma_int_t nb = magma_get_sgehrd_nb( n );
magma_int_t ldda = magma_roundup( n, 32 );
magma_int_t i, nh, iws;
magma_int_t iinfo;
magma_int_t lquery;
*info = 0;
iws = n*nb;
work[0] = magma_smake_lwork( iws );
lquery = (lwork == -1);
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
// Adjust from 1-based indexing
ilo -= 1;
// Quick return if possible
nh = ihi - ilo;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
// If not enough workspace, use unblocked code
if ( lwork < iws ) {
nb = 1;
}
if (nb == 1 || nb > nh) {
// Use unblocked code below
i = ilo;
}
else {
// Use blocked code
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
// GPU workspace is:
// nb*ldda for dwork for slahru
// nb*ldda for dV
// n*ldda for dA
// nb*nb for dT
magmaFloat_ptr dwork;
if (MAGMA_SUCCESS != magma_smalloc( &dwork, 2*nb*ldda + n*ldda + nb*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
float *dV = dwork + nb*ldda;
float *dA = dwork + nb*ldda*2;
float *dT = dwork + nb*ldda*2 + n*ldda;
float *T;
magma_smalloc_cpu( &T, nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// zero first block of V, which is lower triangular
magmablas_slaset( MagmaFull, nb, nb, c_zero, c_zero, dV(0,0), ldda, queue );
// Set elements 0:ILO-1 and IHI-1:N-2 of TAU to zero
for (i = 0; i < ilo; ++i)
tau[i] = c_zero;
for (i = max(0,ihi-1); i < n-1; ++i)
tau[i] = c_zero;
assert( nb % 4 == 0 );
for (i=0; i < nb*nb; i += 4)
T[i] = T[i+1] = T[i+2] = T[i+3] = c_zero;
// Copy the matrix to the GPU
magma_ssetmatrix( n, n-ilo, A(0,ilo), lda, dA(0,0), ldda, queue );
for (i = ilo; i < ihi-1 - nb; i += nb) {
// Reduce columns i:i+nb-1 to Hessenberg form, returning the
// matrices V and T of the block reflector H = I - V*T*V'
// which performs the reduction, and also the matrix Y = A*V*T
// Get the current panel (no need for the 1st iteration)
magma_sgetmatrix( ihi-i, nb,
dA(i,i-ilo), ldda,
A(i,i), lda, queue );
// add 1 to i for 1-based index
magma_slahr2( ihi, i+1, nb,
dA(0,i-ilo), ldda,
dV(0,0), ldda,
A(0,i), lda,
&tau[i], T, nb, work, n, queue );
// Copy T from the CPU to dT on the GPU
magma_ssetmatrix( nb, nb, T, nb, dT(0,0), nb, queue );
magma_slahru( n, ihi, i, nb,
A(0,i), lda,
dA(0,i-ilo), ldda, // dA
dA(i,i-ilo), ldda, // dY, stored over current panel
dV(0,0), ldda,
dT(0,0), dwork, queue );
}
// Copy remainder to host
magma_sgetmatrix( n, n-i,
dA(0,i-ilo), ldda,
A(0,i), lda, queue );
magma_free( dwork );
magma_free_cpu( T );
magma_queue_destroy( queue );
}
// Use unblocked code to reduce the rest of the matrix
// add 1 to i for 1-based index
i += 1;
lapackf77_sgehd2(&n, &i, &ihi, A, &lda, tau, work, &iinfo);
work[0] = magma_smake_lwork( iws );
return *info;
} /* magma_sgehrd2 */
extern "C" magma_int_t
magma_slobpcg( magma_s_sparse_matrix A, magma_s_solver_par *solver_par ) {
#define residualNorms(i,iter) ( residualNorms + (i) + (iter)*n )
#define magmablas_swap(x, y) { pointer = x; x = y; y = pointer; }
#define hresidualNorms(i,iter) (hresidualNorms + (i) + (iter)*n )
#define gramA( m, n) (gramA + (m) + (n)*ldgram)
#define gramB( m, n) (gramB + (m) + (n)*ldgram)
#define gevectors(m, n) (gevectors + (m) + (n)*ldgram)
#define h_gramB( m, n) (h_gramB + (m) + (n)*ldgram)
#define magma_s_bspmv_tuned(m, n, alpha, A, X, beta, AX) { \
magmablas_stranspose( m, n, X, m, blockW, n ); \
magma_s_vector x, ax; \
x.memory_location = Magma_DEV; x.num_rows = m*n; x.nnz = m*n; x.val = blockW; \
ax.memory_location= Magma_DEV; ax.num_rows = m*n; ax.nnz = m*n; ax.val = AX; \
magma_s_spmv(alpha, A, x, beta, ax ); \
magmablas_stranspose( n, m, blockW, n, X, m ); \
}
//**************************************************************
// Memory allocation for the eigenvectors, eigenvalues, and workspace
solver_par->solver = Magma_LOBPCG;
magma_int_t m = A.num_rows;
magma_int_t n =(solver_par->num_eigenvalues);
float *blockX = solver_par->eigenvectors;
float *evalues = solver_par->eigenvalues;
float *dwork, *hwork;
float *blockP, *blockAP, *blockR, *blockAR, *blockAX, *blockW;
float *gramA, *gramB, *gramM;
float *gevectors, *h_gramB;
float *pointer, *origX = blockX;
float *eval_gpu;
magma_int_t lwork = max( 2*n+n*magma_get_dsytrd_nb(n),
1 + 6*3*n + 2* 3*n* 3*n);
magma_smalloc_pinned( &hwork , lwork );
magma_smalloc( &blockAX , m*n );
magma_smalloc( &blockAR , m*n );
magma_smalloc( &blockAP , m*n );
magma_smalloc( &blockR , m*n );
magma_smalloc( &blockP , m*n );
magma_smalloc( &blockW , m*n );
magma_smalloc( &dwork , m*n );
magma_smalloc( &eval_gpu , 3*n );
//**********************************************************+
magma_int_t verbosity = 1;
magma_int_t *iwork, liwork = 15*n+9;
// === Set solver parameters ===
float residualTolerance = solver_par->epsilon;
magma_int_t maxIterations = solver_par->maxiter;
// === Set some constants & defaults ===
float c_one = MAGMA_S_ONE, c_zero = MAGMA_S_ZERO;
float *residualNorms, *condestGhistory, condestG;
float *gevalues;
magma_int_t *activeMask;
// === Check some parameters for possible quick exit ===
solver_par->info = 0;
if (m < 2)
solver_par->info = -1;
else if (n > m)
solver_par->info = -2;
if (solver_par->info != 0) {
magma_xerbla( __func__, -(solver_par->info) );
return solver_par->info;
}
magma_int_t *info = &(solver_par->info); // local info variable;
// === Allocate GPU memory for the residual norms' history ===
magma_smalloc(&residualNorms, (maxIterations+1) * n);
magma_malloc( (void **)&activeMask, (n+1) * sizeof(magma_int_t) );
// === Allocate CPU work space ===
magma_smalloc_cpu(&condestGhistory, maxIterations+1);
magma_smalloc_cpu(&gevalues, 3 * n);
magma_malloc_cpu((void **)&iwork, liwork * sizeof(magma_int_t));
float *hW;
magma_smalloc_pinned(&hW, n*n);
magma_smalloc_pinned(&gevectors, 9*n*n);
magma_smalloc_pinned(&h_gramB , 9*n*n);
// === Allocate GPU workspace ===
magma_smalloc(&gramM, n * n);
magma_smalloc(&gramA, 9 * n * n);
magma_smalloc(&gramB, 9 * n * n);
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork = 1 + 5*(3*n) + 2*(3*n)*(3*n);
magma_smalloc_cpu(&rwork, lrwork);
#endif
// === Set activemask to one ===
for(int k =0; k<n; k++)
iwork[k]=1;
magma_setmatrix(n, 1, sizeof(magma_int_t), iwork, n ,activeMask, n);
magma_int_t gramDim, ldgram = 3*n, ikind = 4;
// === Make the initial vectors orthonormal ===
magma_sgegqr_gpu(ikind, m, n, blockX, m, dwork, hwork, info );
//magma_sorthomgs( m, n, blockX );
magma_s_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX );
// === Compute the Gram matrix = (X, AX) & its eigenstates ===
magma_sgemm(MagmaTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n);
magma_ssyevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, evalues, hW, n, hwork, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork, info );
// === Update X = X * evectors ===
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramM, n, c_zero, blockW, m);
magmablas_swap(blockW, blockX);
// === Update AX = AX * evectors ===
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramM, n, c_zero, blockW, m);
magmablas_swap(blockW, blockAX);
condestGhistory[1] = 7.82;
magma_int_t iterationNumber, cBlockSize, restart = 1, iter;
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync();
tempo1=magma_wtime();
// === Main LOBPCG loop ============================================================
for(iterationNumber = 1; iterationNumber < maxIterations; iterationNumber++)
{
// === compute the residuals (R = Ax - x evalues )
magmablas_slacpy( MagmaUpperLower, m, n, blockAX, m, blockR, m);
/*
for(int i=0; i<n; i++){
magma_saxpy(m, MAGMA_S_MAKE(-evalues[i],0), blockX+i*m, 1, blockR+i*m, 1);
}
*/
#if defined(PRECISION_z) || defined(PRECISION_d)
magma_dsetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n );
#else
magma_ssetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n );
#endif
magma_slobpcg_res( m, n, eval_gpu, blockX, blockR, eval_gpu);
magmablas_snrm2_cols(m, n, blockR, m, residualNorms(0, iterationNumber));
// === remove the residuals corresponding to already converged evectors
magma_scompact(m, n, blockR, m,
residualNorms(0, iterationNumber), residualTolerance,
activeMask, &cBlockSize);
if (cBlockSize == 0)
break;
// === apply a preconditioner P to the active residulas: R_new = P R_old
// === for now set P to be identity (no preconditioner => nothing to be done )
// magmablas_slacpy( MagmaUpperLower, m, cBlockSize, blockR, m, blockW, m);
/*
// === make the preconditioned residuals orthogonal to X
magma_sgemm(MagmaTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockR, m, c_zero, gramB(0,0), ldgram);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_mone, blockX, m, gramB(0,0), ldgram, c_one, blockR, m);
*/
// === make the active preconditioned residuals orthonormal
magma_sgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, info );
//magma_sorthomgs( m, cBlockSize, blockR );
// === compute AR
magma_s_bspmv_tuned(m, cBlockSize, c_one, A, blockR, c_zero, blockAR );
if (!restart) {
// === compact P & AP as well
magma_scompactActive(m, n, blockP, m, activeMask);
magma_scompactActive(m, n, blockAP, m, activeMask);
/*
// === make P orthogonal to X ?
magma_sgemm(MagmaTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockP, m, c_zero, gramB(0,0), ldgram);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_mone, blockX, m, gramB(0,0), ldgram, c_one, blockP, m);
// === make P orthogonal to R ?
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockP, m, c_zero, gramB(0,0), ldgram);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, cBlockSize,
c_mone, blockR, m, gramB(0,0), ldgram, c_one, blockP, m);
*/
// === Make P orthonormal & properly change AP (without multiplication by A)
magma_sgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, info );
//magma_sorthomgs( m, cBlockSize, blockP );
//magma_s_bspmv_tuned(m, cBlockSize, c_one, A, blockP, c_zero, blockAP );
magma_ssetmatrix( cBlockSize, cBlockSize, hwork, cBlockSize, dwork, cBlockSize);
// magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
// m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
// replacement according to Stan
#if defined(PRECISION_s) || defined(PRECISION_d)
magmablas_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
#else
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m,
cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
#endif
}
iter = max(1,iterationNumber-10- (int)(log(1.*cBlockSize)));
float condestGmean = 0.;
for(int i = 0; i<iterationNumber-iter+1; i++)
condestGmean += condestGhistory[i];
condestGmean = condestGmean / (iterationNumber-iter+1);
if (restart)
gramDim = n+cBlockSize;
else
gramDim = n+2*cBlockSize;
/* --- The Raileight-Ritz method for [X R P] -----------------------
[ X R P ]' [AX AR AP] y = evalues [ X R P ]' [ X R P ], i.e.,
GramA GramB
/ X'AX X'AR X'AP \ / X'X X'R X'P \
| R'AX R'AR R'AP | y = evalues | R'X R'R R'P |
\ P'AX P'AR P'AP / \ P'X P'R P'P /
----------------------------------------------------------------- */
// === assemble GramB; first, set it to I
magmablas_slaset(MagmaFull, ldgram, ldgram, c_zero, c_one, gramB, ldgram); // identity
if (!restart) {
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockX, m, c_zero, gramB(n+cBlockSize,0), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockR, m, c_zero, gramB(n+cBlockSize,n), ldgram);
}
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockX, m, c_zero, gramB(n,0), ldgram);
// === get GramB from the GPU to the CPU and compute its eigenvalues only
magma_sgetmatrix(gramDim, gramDim, gramB, ldgram, h_gramB, ldgram);
lapackf77_ssyev("N", "L", &gramDim, h_gramB, &ldgram, gevalues,
hwork, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork,
#endif
info);
// === check stability criteria if we need to restart
condestG = log10( gevalues[gramDim-1]/gevalues[0] ) + 1.;
if ((condestG/condestGmean>2 && condestG>2) || condestG>8) {
// Steepest descent restart for stability
restart=1;
printf("restart at step #%d\n", (int) iterationNumber);
}
// === assemble GramA; first, set it to I
magmablas_slaset(MagmaFull, ldgram, ldgram, c_zero, c_one, gramA, ldgram); // identity
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockAX, m, c_zero, gramA(n,0), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockAR, m, c_zero, gramA(n,n), ldgram);
if (!restart) {
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockAX, m, c_zero,
gramA(n+cBlockSize,0), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAR, m, c_zero,
gramA(n+cBlockSize,n), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAP, m, c_zero,
gramA(n+cBlockSize,n+cBlockSize), ldgram);
}
/*
// === Compute X' AX or just use the eigenvalues below ?
magma_sgemm(MagmaTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero,
gramA(0,0), ldgram);
*/
if (restart==0) {
magma_sgetmatrix(gramDim, gramDim, gramA, ldgram, gevectors, ldgram);
}
else {
gramDim = n+cBlockSize;
magma_sgetmatrix(gramDim, gramDim, gramA, ldgram, gevectors, ldgram);
}
for(int k=0; k<n; k++)
*gevectors(k,k) = MAGMA_S_MAKE(evalues[k], 0);
// === the previous eigensolver destroyed what is in h_gramB => must copy it again
magma_sgetmatrix(gramDim, gramDim, gramB, ldgram, h_gramB, ldgram);
magma_int_t itype = 1;
lapackf77_ssygvd(&itype, "V", "L", &gramDim,
gevectors, &ldgram, h_gramB, &ldgram,
gevalues, hwork, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
for(int k =0; k<n; k++)
evalues[k] = gevalues[k];
// === copy back the result to gramA on the GPU and use it for the updates
magma_ssetmatrix(gramDim, gramDim, gevectors, ldgram, gramA, ldgram);
if (restart == 0) {
// === contribution from P to the new X (in new search direction P)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockP);
// === contribution from R to the new X (in new search direction P)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_one, blockP, m);
// === corresponding contribution from AP to the new AX (in AP)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockAP);
// === corresponding contribution from AR to the new AX (in AP)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_one, blockAP, m);
}
else {
// === contribution from R (only) to the new X
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_zero, blockP, m);
// === corresponding contribution from AR (only) to the new AX
magma_sgemm(MagmaNoTrans, MagmaNoTrans,m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_zero, blockAP, m);
}
// === contribution from old X to the new X + the new search direction P
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramA, ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockX);
//magma_saxpy(m*n, c_one, blockP, 1, blockX, 1);
magma_slobpcg_maxpy( m, n, blockP, blockX );
// === corresponding contribution from old AX to new AX + AP
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramA, ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockAX);
//magma_saxpy(m*n, c_one, blockAP, 1, blockAX, 1);
magma_slobpcg_maxpy( m, n, blockAP, blockAX );
condestGhistory[iterationNumber+1]=condestG;
if (verbosity==1) {
// float res;
// magma_sgetmatrix(1, 1,
// (float*)residualNorms(0, iterationNumber), 1,
// (float*)&res, 1);
//
// printf("Iteration %4d, CBS %4d, Residual: %10.7f\n",
// iterationNumber, cBlockSize, res);
printf("%4d-%2d ", (int) iterationNumber, (int) cBlockSize);
magma_sprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
}
restart = 0;
} // === end for iterationNumber = 1,maxIterations =======================
// fill solver info
magma_device_sync();
tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
solver_par->numiter = iterationNumber;
if( solver_par->numiter < solver_par->maxiter) {
solver_par->info = 0;
} else if( solver_par->init_res > solver_par->final_res )
solver_par->info = -2;
else
solver_par->info = -1;
// =============================================================================
// === postprocessing;
// =============================================================================
// === compute the real AX and corresponding eigenvalues
magma_s_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX );
magma_sgemm(MagmaTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n);
magma_ssyevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, gevalues, dwork, n, hwork, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork, info );
for(int k =0; k<n; k++)
evalues[k] = gevalues[k];
// === update X = X * evectors
magmablas_swap(blockX, dwork);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockX, m);
// === update AX = AX * evectors to compute the final residual
magmablas_swap(blockAX, dwork);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockAX, m);
// === compute R = AX - evalues X
magmablas_slacpy( MagmaUpperLower, m, n, blockAX, m, blockR, m);
for(int i=0; i<n; i++)
magma_saxpy(m, MAGMA_S_MAKE(-evalues[i], 0), blockX+i*m, 1, blockR+i*m, 1);
// === residualNorms[iterationNumber] = || R ||
magmablas_snrm2_cols(m, n, blockR, m, residualNorms(0, iterationNumber));
// === restore blockX if needed
if (blockX != origX)
magmablas_slacpy( MagmaUpperLower, m, n, blockX, m, origX, m);
printf("Eigenvalues:\n");
for(int i =0; i<n; i++)
printf("%e ", evalues[i]);
printf("\n\n");
printf("Final residuals:\n");
magma_sprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
printf("\n\n");
//=== Print residual history in a file for plotting ====
float *hresidualNorms;
magma_smalloc_cpu(&hresidualNorms, (iterationNumber+1) * n);
magma_sgetmatrix(n, iterationNumber,
(float*)residualNorms, n,
(float*)hresidualNorms, n);
printf("Residuals are stored in file residualNorms\n");
printf("Plot the residuals using: myplot \n");
FILE *residuals_file;
residuals_file = fopen("residualNorms", "w");
for(int i =1; i<iterationNumber; i++) {
for(int j = 0; j<n; j++)
fprintf(residuals_file, "%f ", *hresidualNorms(j,i));
fprintf(residuals_file, "\n");
}
fclose(residuals_file);
magma_free_cpu(hresidualNorms);
// === free work space
magma_free( residualNorms );
magma_free_cpu( condestGhistory );
magma_free_cpu( gevalues );
magma_free_cpu( iwork );
magma_free_pinned( hW );
magma_free_pinned( gevectors );
magma_free_pinned( h_gramB );
magma_free( gramM );
magma_free( gramA );
magma_free( gramB );
magma_free( activeMask );
magma_free( blockAX );
magma_free( blockAR );
magma_free( blockAP );
magma_free( blockR );
magma_free( blockP );
magma_free( blockW );
magma_free( dwork );
magma_free( eval_gpu );
magma_free_pinned( hwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
magma_free_cpu( rwork );
#endif
return MAGMA_SUCCESS;
}
/**
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 */
extern "C" magma_int_t
magma_sgehrd(magma_int_t n, magma_int_t ilo, magma_int_t ihi,
float *A, magma_int_t lda,
float *tau,
float *work, magma_int_t lwork,
float *dT,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
Purpose
=======
SGEHRD reduces a REAL general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H . This version
stores the triangular matrices used in the factorization so that they can
be applied directly (i.e., without being recomputed) later. As a result,
the application of Q is much faster.
Arguments
=========
N (input) INTEGER
The order of the matrix A. N >= 0.
ILO (input) INTEGER
IHI (input) INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to SGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, 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).
TAU (output) REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
WORK (workspace/output) REAL array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,N).
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.
dT (output) REAL array on the GPU, dimension NB*N,
where NB is the optimal blocksize. It stores the NB*NB blocks
of the triangular T matrices used in the reduction.
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 (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-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, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
This version stores the T matrices in dT, for later use in magma_sorghr.
===================================================================== */
#define A( i, j ) ( A + (i) + (j)*lda)
#define dA( i, j ) (dA + (i) + (j-ilo)*ldda)
float c_one = MAGMA_S_ONE;
float c_zero = MAGMA_S_ZERO;
magma_int_t nb = magma_get_sgehrd_nb(n);
magma_int_t ldda = n; // assumed in slahru
magma_int_t nh, iws;
magma_int_t iinfo;
magma_int_t ldwork;
magma_int_t lquery;
*info = 0;
iws = n*nb;
work[0] = MAGMA_S_MAKE( iws, 0 );
lquery = lwork == -1;
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
// Adjust from 1-based indexing
ilo -= 1;
// Quick return if possible
nh = ihi - ilo;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
// GPU workspace is:
// nb*ldda for dwork for slahru
// nb*ldda for dV
// n*ldda for dA
float *dwork;
if (MAGMA_SUCCESS != magma_smalloc( &dwork, 2*nb*ldda + n*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
float *dV = dwork + nb*ldda;
float *dA = dwork + nb*ldda*2;
ldwork = n;
magma_int_t i;
float *T, *dTi;
magma_smalloc_cpu( &T, nb*nb );
if ( T == NULL ) {
magma_free( dwork );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// zero first block of V, which is lower triangular
szero_nbxnb_block(nb, dV, ldda);
// Set elements 0:ILO-1 and IHI-1:N-2 of TAU to zero
for(i = 0; i < ilo; ++i)
tau[i] = c_zero;
for(i = max(0,ihi-1); i < n-1; ++i)
tau[i] = c_zero;
for(i=0; i < nb*nb; i += 4)
T[i] = T[i+1] = T[i+2] = T[i+3] = c_zero;
magmablas_slaset( 'F', nb, n, dT, nb );
// If not enough workspace, use unblocked code
if ( lwork < iws ) {
nb = 1;
}
if (nb == 1 || nb > nh) {
// Use unblocked code below
i = ilo;
}
else {
// Use blocked code
// Copy the matrix to the GPU
magma_ssetmatrix( n, n-ilo, A(0,ilo), lda, dA, ldda );
for (i = ilo; i < ihi-1 - nb; i += nb) {
// Reduce columns i:i+nb-1 to Hessenberg form, returning the
// matrices V and T of the block reflector H = I - V*T*V'
// which performs the reduction, and also the matrix Y = A*V*T
// Get the current panel (no need for the 1st iteration)
magma_sgetmatrix( ihi-i, nb,
dA(i,i), ldda,
A (i,i), lda );
// add 1 to i for 1-based index
magma_slahr2( ihi, i+1, nb,
dA(0,i),
dV,
A (0,i), lda,
&tau[i], T, nb, work, ldwork);
// Copy T from the CPU to dT on the GPU
dTi = dT + (i - ilo)*nb;
magma_ssetmatrix( nb, nb, T, nb, dTi, nb );
magma_slahru( n, ihi, i, nb,
A (0,i), lda,
dA(0,i), // dA
dA(i,i), // dY, stored over current panel
dV, dTi, dwork );
}
// Copy remainder to host
magma_sgetmatrix( n, n-i,
dA(0,i), ldda,
A (0,i), lda );
}
// Use unblocked code to reduce the rest of the matrix
// add 1 to i for 1-based index
i += 1;
lapackf77_sgehd2(&n, &i, &ihi, A, &lda, tau, work, &iinfo);
work[0] = MAGMA_S_MAKE( iws, 0 );
magma_free( dwork );
magma_free_cpu( T );
return *info;
} /* magma_sgehrd */
extern "C" magma_int_t
magma_sorgqr_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)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
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
=========
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.
DA (input/output) REAL array A on the GPU, 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.
LDDA (input) INTEGER
The first dimension of the array A. LDDA >= max(1,M).
TAU (input) REAL array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by SGEQRF_GPU.
DT (input/workspace) REAL work space array on the GPU,
dimension (2*MIN(M, N) + (N+31)/32*32 )*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 ].
NB (input) 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.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, lpanel;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork;
float *dV, *dW;
float *work, *panel;
*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;
}
// first kk columns are handled by blocked method.
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min( k, ki+nb );
} else {
kk = 0;
}
// Allocate CPU work space
// n*nb for sorgqr workspace
// (m - kk)*(n - kk) for last block's panel
lwork = n*nb;
lpanel = (m - kk)*(n - kk);
magma_smalloc_cpu( &work, lwork + lpanel );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
panel = work + lwork;
// Allocate work space on GPU
if (MAGMA_SUCCESS != magma_smalloc( &dV, ldda*nb )) {
magma_free_cpu( work );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// dT workspace has:
// 2*min(m,n)*nb for T and R^{-1} matrices from geqrf
// ((n+31)/32*32 )*nb for dW larfb workspace.
lddwork = min(m,n);
dW = dT + 2*lddwork*nb;
cudaStream_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;
magma_sgetmatrix( m_kk, n_kk,
dA(kk, kk), ldda, panel, m_kk );
lapackf77_sorgqr( &m_kk, &n_kk, &k_kk,
panel, &m_kk,
&tau[kk], work, &lwork, &iinfo );
magma_ssetmatrix( m_kk, n_kk,
panel, m_kk, dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_slaset( MagmaUpperLower, kk, n - kk, dA(0, kk), ldda );
}
if (kk > 0) {
// Use blocked code
// stream: copy Aii to V --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min( nb, k-i );
mi = m - i;
// Copy current panel on the GPU from dA to dV
magma_scopymatrix_async( mi, ib,
dA(i,i), ldda,
dV, ldda, stream );
// set panel to identity
magmablas_slaset( MagmaUpperLower, i, ib, dA(0, i), ldda );
magmablas_slaset_identity( mi, ib, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
}
magma_queue_sync( stream );
magmablasSetKernelStream( NULL );
magma_free( dV );
magma_free_cpu( work );
magma_queue_destroy( stream );
return *info;
} /* magma_sorgqr_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing slaset
Code is very similar to testing_slacpy.cpp
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gbytes, gpu_perf, gpu_time, cpu_perf, cpu_time;
float error, work[1];
float c_neg_one = MAGMA_S_NEG_ONE;
float *h_A, *h_R;
magmaFloat_ptr d_A;
float offdiag = MAGMA_S_MAKE( 1.2000, 6.7000 );
float diag = MAGMA_S_MAKE( 3.1415, 2.7183 );
magma_int_t M, N, size, lda, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_uplo_t uplo[] = { MagmaLower, MagmaUpper, MagmaFull };
printf("uplo M N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("=================================================================\n");
for( int iuplo = 0; iuplo < 3; ++iuplo ) {
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
//M += 2; // space for insets
//N += 2;
lda = M;
ldda = ((M+31)/32)*32;
size = lda*N;
if ( uplo[iuplo] == MagmaLower || uplo[iuplo] == MagmaUpper ) {
// save triangle (with diagonal)
// TODO wrong for trapezoid
gbytes = sizeof(float) * 0.5*N*(N+1) / 1e9;
}
else {
// save entire matrix
gbytes = sizeof(float) * 1.*M*N / 1e9;
}
TESTING_MALLOC_CPU( h_A, float, size );
TESTING_MALLOC_CPU( h_R, float, size );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
/* Initialize the matrix */
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < M; ++i ) {
h_A[i + j*lda] = MAGMA_S_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix( M, N, h_A, lda, d_A, 0, ldda, opts.queue );
gpu_time = magma_sync_wtime( 0 );
//magmablas_slaset( uplo[iuplo], M-2, N-2, offdiag, diag, d_A+1+ldda, 0, ldda, opts.queue ); // inset by 1 row & col
magmablas_slaset( uplo[iuplo], M, N, offdiag, diag, d_A, 0, ldda, opts.queue );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//magma_int_t M2 = M-2; // inset by 1 row & col
//magma_int_t N2 = N-2;
//lapackf77_slaset( lapack_uplo_const( uplo[iuplo] ), &M2, &N2, &offdiag, &diag, h_A+1+lda, &lda );
lapackf77_slaset( lapack_uplo_const( uplo[iuplo] ), &M, &N, &offdiag, &diag, h_A, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
if ( opts.verbose ) {
printf( "A= " ); magma_sprint( M, N, h_A, lda );
printf( "dA=" ); magma_sprint_gpu( M, N, d_A, 0, ldda, opts.queue );
}
/* =====================================================================
Check the result
=================================================================== */
magma_sgetmatrix( M, N, d_A, 0, ldda, h_R, lda, opts.queue );
blasf77_saxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &M, &N, h_R, &lda, work);
printf("%5s %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
lapack_uplo_const( uplo[iuplo] ), (int) M, (int) N,
cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000.,
(error == 0. ? "ok" : "failed") );
status += ! (error == 0.);
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
printf( "\n" );
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
SLAHR2 reduces the first NB columns of a real general n-BY-(n-k+1)
matrix A so that elements below the k-th subdiagonal are zero. The
reduction is performed by an orthogonal similarity transformation
Q' * A * Q. The routine returns the matrices V and T which determine
Q as a block reflector I - V*T*V', and also the matrix Y = A * V.
(Note this is different than LAPACK, which computes Y = A * V * T.)
This is an auxiliary routine called by SGEHRD.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
k INTEGER
The offset for the reduction. Elements below the k-th
subdiagonal in the first NB columns are reduced to zero.
K < N.
@param[in]
nb INTEGER
The number of columns to be reduced.
@param[in,out]
A REAL array, dimension (LDA,N-K+1)
On entry, the n-by-(n-k+1) general matrix A.
On exit, the elements on and above the k-th subdiagonal in
the first NB columns are overwritten with the corresponding
elements of the reduced matrix; the elements below the k-th
subdiagonal, with the array TAU, represent the matrix Q as a
product of elementary reflectors. The other columns of A are
unchanged. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau REAL array, dimension (NB)
The scalar factors of the elementary reflectors. See Further
Details.
@param[out]
T REAL array, dimension (LDT,NB)
The upper triangular matrix T.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= NB.
@param[out]
Y REAL array, dimension (LDY,NB)
The n-by-nb matrix Y.
@param[in]
ldy INTEGER
The leading dimension of the array Y. LDY >= N.
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
The matrix Q is represented as a product of nb elementary reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i+k-1) = 0, v(i+k) = 1; v(i+k+1:n) is stored on exit in
A(i+k+1:n,i), and tau in TAU(i).
The elements of the vectors v together form the (n-k+1)-by-nb matrix
V which is needed, with T and Y, to apply the transformation to the
unreduced part of the matrix, using an update of the form:
A := (I - V*T*V') * (A - Y*T*V').
The contents of A on exit are illustrated by the following example
with n = 7, k = 3 and nb = 2:
@verbatim
( a a a a a )
( a a a a a )
( a a a a a )
( h h a a a )
( v1 h a a a )
( v1 v2 a a a )
( v1 v2 a a a )
@endverbatim
where "a" denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
@ingroup magma_sgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_slahr2_m(
magma_int_t n, magma_int_t k, magma_int_t nb,
float *A, magma_int_t lda,
float *tau,
float *T, magma_int_t ldt,
float *Y, magma_int_t ldy,
struct sgehrd_data *data )
{
#define A( i, j ) ( A + (i) + (j)*lda)
#define Y( i, j ) ( Y + (i) + (j)*ldy)
#define T( i, j ) ( T + (i) + (j)*ldt)
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
float tmp;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t ione = 1;
magma_int_t d, dki1, dn, nblocks, gblock, lblock, lgid;
magma_int_t n_k_i_1, n_k;
float scale;
magma_int_t i;
float ei = MAGMA_S_ZERO;
magma_int_t info_data = 0;
magma_int_t *info = &info_data;
if (n < 0) {
*info = -1;
} else if (k < 0 || k >= n) {
*info = -2;
} else if (nb < 1 || nb > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldt < nb) {
*info = -8;
} else if (ldy < max(1,n)) {
*info = -10;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
// adjust from 1-based indexing
k -= 1;
// Function Body
if (n <= 1)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// zero out current top block of V on all GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablas_slaset( MagmaFull, nb, nb, c_zero, c_zero, dV(d,k,0), ldv, data->queues[d] );
}
// set all Y=0
lapackf77_slaset( "Full", &n, &nb, &c_zero, &c_zero, Y, &ldy );
for (i = 0; i < nb; ++i) {
n_k_i_1 = n - k - i - 1;
n_k = n - k;
if (i > 0) {
// Finish applying I - V * T * V' on right
tmp = MAGMA_S_NEGATE( tau[i-1] );
blasf77_saxpy( &n_k, &tmp, Y(k,i-1), &ione, A(k,i), &ione );
// Apply I - V * T' * V' to this column (call it b) from the
// left, using the last column of T as workspace, w.
//
// Let V = ( V1 ) and b = ( b1 ) (first i-1 rows)
// ( V2 ) ( b2 )
// where V1 is unit lower triangular
// w := b1 = A(k+1:k+i, i)
blasf77_scopy( &i,
A(k+1,i), &ione,
T(0,nb-1), &ione );
// w := V1' * b1 = VA(k+1:k+i, 0:i-1)' * w
blasf77_strmv( "Lower", "Conj", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// w := w + V2'*b2 = w + VA(k+i+1:n-1, 0:i-1)' * A(k+i+1:n-1, i)
blasf77_sgemv( "Conj", &n_k_i_1, &i,
&c_one, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_one, T(0,nb-1), &ione );
// w := T'*w = T(0:i-1, 0:i-1)' * w
blasf77_strmv( "Upper", "Conj", "Non-unit", &i,
T(0,0), &ldt,
T(0,nb-1), &ione );
// b2 := b2 - V2*w = A(k+i+1:n-1, i) - VA(k+i+1:n-1, 0:i-1) * w
blasf77_sgemv( "No trans", &n_k_i_1, &i,
&c_neg_one, A(k+i+1,0), &lda,
T(0,nb-1), &ione,
&c_one, A(k+i+1,i), &ione );
// w := V1*w = VA(k+1:k+i, 0:i-1) * w
blasf77_strmv( "Lower", "No trans", "Unit", &i,
A(k+1,0), &lda,
T(0,nb-1), &ione );
// b1 := b1 - w = A(k+1:k+i-1, i) - w
blasf77_saxpy( &i,
&c_neg_one, T(0,nb-1), &ione,
A(k+1,i), &ione );
// Restore diagonal element, saved below during previous iteration
*A(k+i,i-1) = ei;
}
// Generate the elementary reflector H(i) to annihilate A(k+i+1:n-1,i)
lapackf77_slarfg( &n_k_i_1,
A(k+i+1,i),
A(k+i+2,i), &ione, &tau[i] );
// Save diagonal element and set to one, to simplify multiplying by V
ei = *A(k+i+1,i);
*A(k+i+1,i) = c_one;
// compute yi = A vi = sum_g A{d} vi{d}
nblocks = (n-1) / nb / ngpu + 1;
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
// dV(k+i+1:n-1, i) = VA(k+i:n, i)
magma_ssetvector_async( n_k_i_1,
A(k+i+1,i), 1,
dV(d, k+i+1, i), 1, data->queues[d] );
// copy column of dV -> dVd, using block cyclic distribution.
// This assumes V and Vd have been padded so that
// a 2D matrix copy doesn't access them out-of-bounds
gblock = k / nb;
lblock = gblock / ngpu;
lgid = gblock % ngpu;
if ( d < lgid ) {
lblock += 1;
}
// treat V as (nb*ngpu) x nblock matrix, and Vd as nb x nblock matrix
magmablas_slacpy( MagmaFull, nb, nblocks-lblock,
dV (d, d*nb + lblock*nb*ngpu, i), nb*ngpu,
dVd(d, 0 + lblock*nb, i), nb, data->queues[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
// dY(k:n, i) = dA(k:n, k+i+1:n) * dV(k+i+1:n, i)
// skip if matrix is empty
// each GPU copies to different temporary vector in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_sgemv( MagmaNoTrans, n-k, dn-dki1,
c_one, dA (d, k, dki1), ldda,
dVd(d, dki1, i), 1,
c_zero, dY (d, k, i), 1, data->queues[d] );
// copy vector to host, storing in column nb+d of Y
// as temporary space (Y has >= nb+ngpu columns)
magma_sgetvector_async( n-k,
dY(d, k, i), 1,
Y(k, nb+d), 1, data->queues[d] );
}
}
// while GPU is doing above Ag*v...
// Compute T(0:i,i) = [ -tau T V' vi ]
// [ tau ]
// T(0:i-1, i) = -tau VA(k+i+1:n-1, 0:i-1)' VA(k+i+1:n-1, i)
scale = MAGMA_S_NEGATE( tau[i] );
blasf77_sgemv( "Conj", &n_k_i_1, &i,
&scale, A(k+i+1,0), &lda,
A(k+i+1,i), &ione,
&c_zero, T(0,i), &ione );
// T(0:i-1, i) = T(0:i-1, 0:i-1) * T(0:i-1, i)
blasf77_strmv( "Upper", "No trans", "Non-unit", &i,
T(0,0), &ldt,
T(0,i), &ione );
*T(i,i) = tau[i];
// apply reflectors to next column, A(i+1), on right only.
// one axpy will be required to finish this, in the next iteration above
if ( i > 0 && i+1 < nb ) {
// Update next column, A(k:n,i+1), applying Q on right.
// One axpy will be required to finish this, in the next iteration
// above, after yi is computed.
// This updates one more row than LAPACK does (row k),
// making block above panel an even multiple of nb.
// Use last column of T as workspace, w.
magma_int_t i1 = i+1;
// If real, conjugate row of V, and undo afterwards
#ifdef COMPLEX
lapackf77_slacgv( &i1, A(k+i1,0), &lda );
#endif
// w = T(0:i, 0:i+1) * VA(k+i+1, 0:i+1)'
// T is now rectangular, so we use gemv instead of trmv as in lapack.
blasf77_sgemv( "No trans", &i, &i1,
&c_one, T(0,0), &ldt,
A(k+i1,0), &lda,
&c_zero, T(0,nb-1), &ione );
#ifdef COMPLEX
lapackf77_slacgv( &i1, A(k+i1,0), &lda );
#endif
// A(k:n, i+1) -= Y(k:n, 0:i) * w
blasf77_sgemv( "No trans", &n_k, &i,
&c_neg_one, Y(k,0), &ldy,
T(0,nb-1), &ione,
&c_one, A(k,i1), &ione );
}
// yi = sum_g yi{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( data->queues[d] );
magma_indices_1D_bcyclic( nb, ngpu, d, k+i+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// yi = yi + yi{d}
blasf77_saxpy( &n_k, &c_one, Y(k,nb+d), &ione, Y(k,i), &ione );
}
}
}
// Restore diagonal element
*A(k+nb,nb-1) = ei;
// compute Y = Am V = sum_g Am{d} V{d} --- top part, Y(0:k-1,:)
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
// dY(0:k, :) = dA(0:k, k+i+1:n-1) * dV(k+i+1:n-1, :)
// skip if matrix is empty
// each GPU copies to different temporary block in Y,
// which are summed in separate loop below
if ( dn-dki1 > 0 ) {
magma_sgemm( MagmaNoTrans, MagmaNoTrans, k, nb, dn-dki1,
c_one, dA (d, 0, dki1), ldda,
dVd(d, dki1, 0), ldvd,
c_zero, dY (d, 0, 0), ldda, data->queues[d] );
// copy result to host, storing in columns [nb + nb*d : nb + nb*(d+1)] of Y
// as temporary space (Y has nb + nb*ngpu columns)
magma_sgetmatrix_async( k, nb,
dY(d, 0, 0), ldda,
Y(0,nb+nb*d), ldy, data->queues[d] );
}
}
// Y = sum_g Y{d}
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_queue_sync( 0 );
magma_indices_1D_bcyclic( nb, ngpu, d, k+1, n, &dki1, &dn );
if ( dn-dki1 > 0 ) {
// Y = Y + Am V
for( i = 0; i < nb; ++i ) {
blasf77_saxpy( &k, &c_one, Y(0,nb+nb*d+i), &ione, Y(0,i), &ione );
}
}
}
// copy Y and T matrices to GPUs
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magma_ssetmatrix_async( n, nb, Y, ldy, dY(d, 0, 0), ldda, data->queues[d] );
magma_ssetmatrix_async( nb, nb, T, nb, dTi(d), nb, data->queues[d] );
}
magma_setdevice( orig_dev );
return *info;
} /* magma_slahr2 */
extern "C" magma_int_t
magma_sidr_strms(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_IDRMERGE;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->init_res = 0.0;
solver_par->final_res = 0.0;
solver_par->iter_res = 0.0;
solver_par->runtime = 0.0;
// constants
const float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_n_one = MAGMA_S_NEG_ONE;
// internal user options
const magma_int_t smoothing = 1; // 0 = disable, 1 = enable
const float angle = 0.7; // [0-1]
// local variables
magma_int_t iseed[4] = {0, 0, 0, 1};
magma_int_t dof;
magma_int_t s;
magma_int_t distr;
magma_int_t k, i, sk;
magma_int_t innerflag;
magma_int_t ldd;
magma_int_t q;
float residual;
float nrm;
float nrmb;
float nrmr;
float nrmt;
float rho;
float om;
float gamma;
// matrices and vectors
magma_s_matrix dxs = {Magma_CSR};
magma_s_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_s_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_s_matrix dG = {Magma_CSR}, dGcol = {Magma_CSR};
magma_s_matrix dU = {Magma_CSR};
magma_s_matrix dM = {Magma_CSR};
magma_s_matrix df = {Magma_CSR};
magma_s_matrix dt = {Magma_CSR}, dtt = {Magma_CSR};
magma_s_matrix dc = {Magma_CSR};
magma_s_matrix dv = {Magma_CSR};
magma_s_matrix dskp = {Magma_CSR};
magma_s_matrix dalpha = {Magma_CSR};
magma_s_matrix dbeta = {Magma_CSR};
float *hMdiag = NULL;
float *hskp = NULL;
float *halpha = NULL;
float *hbeta = NULL;
float *d1 = NULL, *d2 = NULL;
// queue variables
const magma_int_t nqueues = 3; // number of queues
magma_queue_t queues[nqueues];
// chronometry
real_Double_t tempo1, tempo2;
// create additional queues
queues[0] = queue;
for ( q = 1; q < nqueues; q++ ) {
magma_queue_create( queue->device(), &(queues[q]) );
}
// initial s space
// TODO: add option for 's' (shadow space number)
// Hack: uses '--restart' option as the shadow space number.
// This is not a good idea because the default value of restart option is used to detect
// if the user provided a custom restart. This means that if the default restart value
// is changed then the code will think it was the user (unless the default value is
// also updated in the 'if' statement below.
s = 1;
if ( solver_par->restart != 50 ) {
if ( solver_par->restart > A.num_cols ) {
s = A.num_cols;
} else {
s = solver_par->restart;
}
}
solver_par->restart = s;
// set max iterations
solver_par->maxiter = min( 2 * A.num_cols, solver_par->maxiter );
// check if matrix A is square
if ( A.num_rows != A.num_cols ) {
//printf("Matrix A is not square.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
goto cleanup;
}
// |b|
nrmb = magma_snrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_sscal( x->num_rows, MAGMA_S_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// t = 0
// make t twice as large to contain both, dt and dr
ldd = magma_roundup( b.num_rows, 32 );
CHECK( magma_svinit( &dt, Magma_DEV, ldd, 2, c_zero, queue ));
dt.num_rows = b.num_rows;
dt.num_cols = 1;
dt.nnz = dt.num_rows;
// redirect the dr.dval to the second part of dt
CHECK( magma_svinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
magma_free( dr.dval );
dr.dval = dt.dval + ldd;
// r = b - A x
CHECK( magma_sresidualvec( A, b, *x, &dr, &nrmr, queue ));
// |r|
solver_par->init_res = nrmr;
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t)nrmr;
}
// check if initial is guess good enough
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
// P = randn(n, s)
// P = ortho(P)
//---------------------------------------
// P = 0.0
CHECK( magma_svinit( &dP, Magma_CPU, A.num_cols, s, c_zero, queue ));
// P = randn(n, s)
distr = 3; // 1 = unif (0,1), 2 = unif (-1,1), 3 = normal (0,1)
dof = dP.num_rows * dP.num_cols;
lapackf77_slarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_smtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_smfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_sqr(P1), QR factorization
CHECK( magma_sqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_snrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_sscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_smtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_smfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_smalloc_pinned( &hskp, 5 ));
CHECK( magma_svinit( &dskp, Magma_DEV, 4, 1, c_zero, queue ));
CHECK( magma_smalloc_pinned( &halpha, s ));
CHECK( magma_svinit( &dalpha, Magma_DEV, s, 1, c_zero, queue ));
CHECK( magma_smalloc_pinned( &hbeta, s ));
CHECK( magma_svinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// workspace for merged dot product
CHECK( magma_smalloc( &d1, max(2, s) * b.num_rows ));
CHECK( magma_smalloc( &d2, max(2, s) * b.num_rows ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_smtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// tt = 0
// make tt twice as large to contain both, dtt and drs
ldd = magma_roundup( b.num_rows, 32 );
CHECK( magma_svinit( &dtt, Magma_DEV, ldd, 2, c_zero, queue ));
dtt.num_rows = dr.num_rows;
dtt.num_cols = 1;
dtt.nnz = dtt.num_rows;
// redirect the drs.dval to the second part of dtt
CHECK( magma_svinit( &drs, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
magma_free( drs.dval );
drs.dval = dtt.dval + ldd;
// set smoothing residual vector
magma_scopyvector( dr.num_rows, dr.dval, 1, drs.dval, 1, queue );
}
// G(n,s) = 0
if ( s > 1 ) {
ldd = magma_roundup( A.num_rows, 32 );
CHECK( magma_svinit( &dG, Magma_DEV, ldd, s, c_zero, queue ));
dG.num_rows = A.num_rows;
} else {
CHECK( magma_svinit( &dG, Magma_DEV, A.num_rows, s, c_zero, queue ));
}
// dGcol represents a single column of dG, array pointer is set inside loop
CHECK( magma_svinit( &dGcol, Magma_DEV, dG.num_rows, 1, c_zero, queue ));
magma_free( dGcol.dval );
// U(n,s) = 0
if ( s > 1 ) {
ldd = magma_roundup( A.num_cols, 32 );
CHECK( magma_svinit( &dU, Magma_DEV, ldd, s, c_zero, queue ));
dU.num_rows = A.num_cols;
} else {
CHECK( magma_svinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
}
// M(s,s) = I
CHECK( magma_svinit( &dM, Magma_DEV, s, s, c_zero, queue ));
CHECK( magma_smalloc_pinned( &hMdiag, s ));
magmablas_slaset( MagmaFull, dM.num_rows, dM.num_cols, c_zero, c_one, dM.dval, dM.ld, queue );
// f = 0
CHECK( magma_svinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// c = 0
CHECK( magma_svinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = r
CHECK( magma_smtransfer( dr, &dv, Magma_DEV, Magma_DEV, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
cudaProfilerStart();
om = MAGMA_S_ONE;
gamma = MAGMA_S_ZERO;
innerflag = 0;
// new RHS for small systems
// f = P' r
// Q1
magma_sgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_sgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_scopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
// start iteration
do
{
solver_par->numiter++;
// shadow space loop
for ( k = 0; k < s; ++k ) {
sk = s - k;
dGcol.dval = dG.dval + k * dG.ld;
// v = r - G(:,k:s) c(k:s)
// Q1
magmablas_sgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, dGcol.dval, dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queues[1] );
// U(:,k) = om * v + U(:,k:s) c(k:s)
// Q1
magmablas_sgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queues[1] );
// G(:,k) = A U(:,k)
// Q1
CHECK( magma_s_spmv( c_one, A, dv, c_zero, dGcol, queues[1] ));
solver_par->spmv_count++;
// bi-orthogonalize the new basis vectors
for ( i = 0; i < k; ++i ) {
// alpha = P(:,i)' G(:,k)
// Q1
halpha[i] = magma_sdot( dP.num_rows, &dP.dval[i*dP.ld], 1, dGcol.dval, 1, queues[1] );
// implicit sync Q1 --> alpha = P(:,i)' G(:,k)
// alpha = alpha / M(i,i)
halpha[i] = halpha[i] / hMdiag[i];
// G(:,k) = G(:,k) - alpha * G(:,i)
// Q1
magma_saxpy( dG.num_rows, -halpha[i], &dG.dval[i*dG.ld], 1, dGcol.dval, 1, queues[1] );
}
// sync Q1 --> G(:,k) = G(:,k) - alpha * G(:,i), skp[4] = f(k)
magma_queue_sync( queues[1] );
// new column of M = P'G, first k-1 entries are zero
// M(k:s,k) = P(:,k:s)' G(:,k)
// Q2
magma_sgemvmdot_shfl( dP.num_rows, sk, &dP.dval[k*dP.ld], dGcol.dval, d1, d2, &dM.dval[k*dM.ld+k], queues[2] );
// non-first s iteration
if ( k > 0 ) {
// alpha = dalpha
// Q0
magma_ssetvector_async( k, halpha, 1, dalpha.dval, 1, queues[0] );
// U update outside of loop using GEMV
// U(:,k) = U(:,k) - U(:,1:k) * alpha(1:k)
// Q0
magmablas_sgemv( MagmaNoTrans, dU.num_rows, k, c_n_one, dU.dval, dU.ld, dalpha.dval, 1, c_one, dv.dval, 1, queues[0] );
}
// Mdiag(k) = M(k,k)
// Q2
magma_sgetvector( 1, &dM.dval[k*dM.ld+k], 1, &hMdiag[k], 1, queues[2] );
// implicit sync Q2 --> Mdiag(k) = M(k,k)
// U(:,k) = v
// Q0
magma_scopyvector_async( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queues[0] );
// check M(k,k) == 0
if ( MAGMA_S_EQUAL(hMdiag[k], MAGMA_S_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
hbeta[k] = hskp[4] / hMdiag[k];
// check for nan
if ( magma_s_isnan( hbeta[k] ) || magma_s_isinf( hbeta[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
// Q2
magma_saxpy( dr.num_rows, -hbeta[k], dGcol.dval, 1, dr.dval, 1, queues[2] );
// non-last s iteration
if ( (k + 1) < s ) {
// f(k+1:s) = f(k+1:s) - beta * M(k+1:s,k)
// Q1
magma_saxpy( sk-1, -hbeta[k], &dM.dval[k*dM.ld+(k+1)], 1, &df.dval[k+1], 1, queues[1] );
// c(k+1:s) = f(k+1:s)
// Q1
magma_scopyvector_async( sk-1, &df.dval[k+1], 1, &dc.dval[k+1], 1, queues[1] );
// c(k+1:s) = M(k+1:s,k+1:s) \ f(k+1:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, sk-1, &dM.dval[(k+1)*dM.ld+(k+1)], dM.ld, &dc.dval[k+1], 1, queues[1] );
// skp[4] = f(k+1)
// Q1
magma_sgetvector_async( 1, &df.dval[k+1], 1, &hskp[4], 1, queues[1] );
}
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
// Q2
nrmr = magma_snrm2( dr.num_rows, dr.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
// Q2
magma_sidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queues[2] );
// x = x + beta * U(:,k)
// Q0
magma_saxpy( x->num_rows, hbeta[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queues[0] );
// t't
// t'rs
// Q2
CHECK( magma_sgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queues[2] ));
// skp[2-3] = dskp[2-3]
// Q2
magma_sgetvector( 2, &dskp.dval[2], 1, &hskp[2], 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// gamma = (t' * rs) / (t' * t)
gamma = hskp[3] / hskp[2];
// rs = rs - gamma * t
// Q1
magma_saxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queues[1] );
// xs = xs - gamma * (xs - x)
// Q0
magma_sidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queues[0] );
// |rs|
// Q1
nrmr = magma_snrm2( drs.num_rows, drs.dval, 1, queues[1] );
// implicit sync Q0 --> |r|
//---------------------------------------
}
// v = r
// Q1
magma_scopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[1] );
// last s iteration
if ( (k + 1) == s ) {
// t = A r
// Q2
CHECK( magma_s_spmv( c_one, A, dr, c_zero, dt, queues[2] ));
solver_par->spmv_count++;
// t't
// t'r
// Q2
CHECK( magma_sgemvmdot_shfl( dt.ld, 2, dt.dval, dt.dval, d1, d2, dskp.dval, queues[2] ));
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter) % solver_par->verbose == 0 ) {
solver_par->res_vec[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence or iteration limit
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
s = k + 1; // for the x-update outside the loop
innerflag = 2;
info = MAGMA_SUCCESS;
break;
}
}
// smoothing disabled
if ( smoothing <= 0 && innerflag != 1 ) {
// dbeta(1:s) = beta(1:s)
// Q0
magma_ssetvector_async( s, hbeta, 1, dbeta.dval, 1, queues[0] );
// x = x + U(:,1:s) * beta(1:s)
// Q0
magmablas_sgemv( MagmaNoTrans, dU.num_rows, s, c_one, dU.dval, dU.ld, dbeta.dval, 1, c_one, x->dval, 1, queues[0] );
}
// check convergence or iteration limit or invalid result of inner loop
if ( innerflag > 0 ) {
break;
}
// computation of a new omega
//---------------------------------------
// skp[0-2] = dskp[0-2]
// Q2
magma_sgetvector( 2, dskp.dval, 1, hskp, 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// |t|
nrmt = magma_ssqrt( MAGMA_S_REAL(hskp[0]) );
// rho = abs((t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_S_REAL(hskp[1]) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = hskp[1] / hskp[0];
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_S_EQUAL(om, MAGMA_S_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// sync Q1 --> v = r
magma_queue_sync( queues[1] );
// r = r - om * t
// Q2
magma_saxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queues[2] );
// x = x + om * v
// Q0
magma_saxpy( x->num_rows, om, dv.dval, 1, x->dval, 1, queues[0] );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
// Q2
nrmr = magma_snrm2( dr.num_rows, dr.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// v = r
// Q0
magma_scopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[0] );
// new RHS for small systems
// f = P' r
// Q1
magma_sgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_sgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_scopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
// Q2
magma_sidr_smoothing_1( drs.num_rows, drs.num_cols, drs.dval, dr.dval, dtt.dval, queues[2] );
// t't
// t'rs
// Q2
CHECK( magma_sgemvmdot_shfl( dt.ld, 2, dtt.dval, dtt.dval, d1, d2, &dskp.dval[2], queues[2] ));
// skp[2-3] = dskp[2-3]
// Q2
magma_sgetvector( 2, &dskp.dval[2], 1, &hskp[2], 1, queues[2] );
// implicit sync Q2 --> skp = dskp
// gamma = (t' * rs) / (t' * t)
gamma = hskp[3] / hskp[2];
// rs = rs - gamma * (rs - r)
// Q2
magma_saxpy( drs.num_rows, -gamma, dtt.dval, 1, drs.dval, 1, queues[2] );
// xs = xs - gamma * (xs - x)
// Q0
magma_sidr_smoothing_2( dxs.num_rows, dxs.num_cols, -gamma, x->dval, dxs.dval, queues[0] );
// v = r
// Q0
magma_scopyvector_async( dr.num_rows, dr.dval, 1, dv.dval, 1, queues[0] );
// new RHS for small systems
// f = P' r
// Q1
magma_sgemvmdot_shfl( dP.num_rows, dP.num_cols, dP.dval, dr.dval, d1, d2, df.dval, queues[1] );
// skp[4] = f(k)
// Q1
magma_sgetvector_async( 1, df.dval, 1, &hskp[4], 1, queues[1] );
// c(k:s) = f(k:s)
// Q1
magma_scopyvector_async( s, df.dval, 1, dc.dval, 1, queues[1] );
// |rs|
// Q2
nrmr = magma_snrm2( drs.num_rows, drs.dval, 1, queues[2] );
// implicit sync Q2 --> |r|
// c(k:s) = M(k:s,k:s) \ f(k:s)
// Q1
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, s, dM.dval, dM.ld, dc.dval, 1, queues[1] );
//---------------------------------------
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
magma_queue_sync( queue );
if ( (solver_par->numiter) % solver_par->verbose == 0 ) {
solver_par->res_vec[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence or iteration limit
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
break;
}
// sync Q0 --> v = r
magma_queue_sync( queues[0] );
}
while ( solver_par->numiter + 1 <= solver_par->maxiter );
// sync all queues
for ( q = 0; q < nqueues; q++ ) {
magma_queue_sync( queues[q] );
}
// smoothing enabled
if ( smoothing > 0 ) {
// x = xs
magma_scopyvector_async( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_scopyvector_async( dr.num_rows, drs.dval, 1, dr.dval, 1, queue );
}
cudaProfilerStop();
// get last iteration timing
tempo2 = magma_sync_wtime( queue );
magma_queue_sync( queue );
solver_par->runtime = (real_Double_t)tempo2 - tempo1;
//--------------STOP TIME----------------
// get final stats
solver_par->iter_res = nrmr;
CHECK( magma_sresidualvec( A, b, *x, &dr, &residual, queue ));
solver_par->final_res = residual;
// set solver conclusion
if ( info != MAGMA_SUCCESS && info != MAGMA_DIVERGENCE ) {
if ( solver_par->init_res > solver_par->final_res ) {
info = MAGMA_SLOW_CONVERGENCE;
}
}
cleanup:
// free resources
// sync all queues, destory additional queues
magma_queue_sync( queues[0] );
for ( q = 1; q < nqueues; q++ ) {
magma_queue_sync( queues[q] );
magma_queue_destroy( queues[q] );
}
// smoothing enabled
if ( smoothing > 0 ) {
drs.dval = NULL; // needed because its pointer is redirected to dtt
magma_smfree( &dxs, queue );
magma_smfree( &drs, queue );
magma_smfree( &dtt, queue );
}
dr.dval = NULL; // needed because its pointer is redirected to dt
dGcol.dval = NULL; // needed because its pointer is redirected to dG
magma_smfree( &dr, queue );
magma_smfree( &dP, queue );
magma_smfree( &dP1, queue );
magma_smfree( &dG, queue );
magma_smfree( &dGcol, queue );
magma_smfree( &dU, queue );
magma_smfree( &dM, queue );
magma_smfree( &df, queue );
magma_smfree( &dt, queue );
magma_smfree( &dc, queue );
magma_smfree( &dv, queue );
magma_smfree( &dskp, queue );
magma_smfree( &dalpha, queue );
magma_smfree( &dbeta, queue );
magma_free_pinned( hMdiag );
magma_free_pinned( hskp );
magma_free_pinned( halpha );
magma_free_pinned( hbeta );
magma_free( d1 );
magma_free( d2 );
solver_par->info = info;
return info;
/* magma_sidr_strms */
}
extern "C" magma_int_t
magma_sgehrd(
magma_int_t n, magma_int_t ilo, magma_int_t ihi,
float *a, magma_int_t lda,
float *tau,
float *work, magma_int_t lwork,
magmaFloat_ptr dT, size_t dT_offset,
magma_queue_t queue,
magma_int_t *info)
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
SGEHRD reduces a REAL general matrix A to upper Hessenberg form H by
an orthogonal similarity transformation: Q' * A * Q = H . This version
stores the triangular matrices used in the factorization so that they can
be applied directly (i.e., without being recomputed) later. As a result,
the application of Q is much faster.
Arguments
=========
N (input) INTEGER
The order of the matrix A. N >= 0.
ILO (input) INTEGER
IHI (input) INTEGER
It is assumed that A is already upper triangular in rows
and columns 1:ILO-1 and IHI+1:N. ILO and IHI are normally
set by a previous call to SGEBAL; otherwise they should be
set to 1 and N respectively. See Further Details.
1 <= ILO <= IHI <= N, if N > 0; ILO=1 and IHI=0, if N=0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the N-by-N general matrix to be reduced.
On exit, the upper triangle and the first subdiagonal of A
are overwritten with the upper Hessenberg matrix H, 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).
TAU (output) REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details). Elements 1:ILO-1 and IHI:N-1 of TAU are set to
zero.
WORK (workspace/output) REAL array, dimension (LWORK)
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK. LWORK >= max(1,N).
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.
dT (output) REAL array on the GPU, dimension N*NB,
where NB is the optimal blocksize. It stores the NB*NB blocks
of the triangular T matrices, used the the reduction.
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 (ihi-ilo) elementary
reflectors
Q = H(ilo) H(ilo+1) . . . H(ihi-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, v(i+1) = 1 and v(ihi+1:n) = 0; v(i+2:ihi) is stored on
exit in A(i+2:ihi,i), and tau in TAU(i).
The contents of A are illustrated by the following example, with
n = 7, ilo = 2 and ihi = 6:
on entry, on exit,
( a a a a a a a ) ( a a h h h h a )
( a a a a a a ) ( a h h h h a )
( a a a a a a ) ( h h h h h h )
( a a a a a a ) ( v2 h h h h h )
( a a a a a a ) ( v2 v3 h h h h )
( a a a a a a ) ( v2 v3 v4 h h h )
( a ) ( a )
where a denotes an element of the original matrix A, h denotes a
modified element of the upper Hessenberg matrix H, and vi denotes an
element of the vector defining H(i).
This implementation follows the hybrid algorithm and notations described in
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
===================================================================== */
float c_one = MAGMA_S_ONE;
float c_zero = MAGMA_S_ZERO;
magma_int_t nb = magma_get_sgehrd_nb(n);
magma_int_t N = n, ldda = n;
magma_int_t ib;
magma_int_t nh, iws;
magma_int_t nbmin, iinfo;
magma_int_t ldwork;
magma_int_t lquery;
--tau;
*info = 0;
work[0] = MAGMA_S_MAKE( n * nb, 0 );
lquery = lwork == -1;
if (n < 0) {
*info = -1;
} else if (ilo < 1 || ilo > max(1,n)) {
*info = -2;
} else if (ihi < min(ilo,n) || ihi > n) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (lwork < max(1,n) && ! lquery) {
*info = -8;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
nh = ihi - ilo + 1;
if (nh <= 1) {
work[0] = c_one;
return *info;
}
magmaFloat_ptr da;
size_t da_offset = 0;
if (MAGMA_SUCCESS != magma_smalloc( &da, (N*ldda + 2*N*nb + nb*nb) )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmaFloat_ptr d_A = da;
size_t d_A_offset = da_offset;
//float *d_work = da + (N+nb)*ldda;
magmaFloat_ptr d_work = da;
size_t d_work_offset = da_offset+(N+nb)*ldda;
magma_int_t i__;
float *t;
magma_smalloc_cpu( &t, nb*nb );
if ( t == NULL ) {
magma_free( da );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magmaFloat_ptr d_t;
d_t = d_work;
size_t d_t_offset = d_work_offset+nb*ldda;
magmablas_slaset( MagmaFull, nb, nb, c_zero, c_zero, d_A, d_A_offset+N*ldda, ldda, queue );
/* Set elements 1:ILO-1 and IHI:N-1 of TAU to zero */
for (i__ = 1; i__ < ilo; ++i__)
tau[i__] = c_zero;
for (i__ = max(1,ihi); i__ < n; ++i__)
tau[i__] = c_zero;
for(i__=0; i__< nb*nb; i__+=4)
t[i__] = t[i__+1] = t[i__+2] = t[i__+3] = c_zero;
nbmin = 2;
iws = 1;
if (nb > 1 && nb < nh) {
/* Determine when to cross over from blocked to unblocked code
(last block is always handled by unblocked code) */
if (nb < nh) {
/* Determine if workspace is large enough for blocked code */
iws = n * nb;
if (lwork < iws) {
/* Not enough workspace to use optimal NB: determine the
minimum value of NB, and reduce NB or force use of
unblocked code */
nbmin = nb;
if (lwork >= n * nbmin)
nb = lwork / n;
else
nb = 1;
}
}
}
ldwork = n;
if (nb < nbmin || nb >= nh) {
/* Use unblocked code below */
i__ = ilo;
} else {
/* Use blocked code */
/* Copy the matrix to the GPU */
magma_ssetmatrix( N, N-ilo+1, a+(ilo-1)*(lda), lda, d_A, d_A_offset, ldda, queue );
for (i__ = ilo; i__ < ihi - nb; i__ += nb) {
/* Computing MIN */
ib = min(nb, ihi - i__);
/* Reduce columns i:i+ib-1 to Hessenberg form, returning the
matrices V and T of the block reflector H = I - V*T*V'
which performs the reduction, and also the matrix Y = A*V*T */
/* Get the current panel (no need for the 1st iteration) */
magma_sgetmatrix( ihi-i__+1, ib,
d_A, (d_A_offset + (i__ - ilo)*ldda + i__ - 1), ldda,
a + (i__ - 1 )*lda + i__ - 1, lda, queue );
magma_slahr2(ihi, i__, ib,
d_A, d_A_offset +(i__ - ilo)*ldda, ldda,
d_A, d_A_offset + N*ldda + 1, ldda,
a + (i__ - 1 )*(lda) , lda,
&tau[i__], t, nb, work, ldwork, queue);
/* Copy T from the CPU to D_T on the GPU */
//d_t = dT + (i__ - ilo)*nb;
d_t = dT;
d_t_offset = dT_offset + (i__ - ilo)*nb;
magma_ssetmatrix( nb, nb, t, nb, d_t, d_t_offset, nb, queue );
magma_slahru(n, ihi, i__ - 1, ib,
a + (i__ - 1 )*(lda), lda,
d_A, d_A_offset + (i__ - ilo)*ldda, ldda,
d_A, d_A_offset + (i__ - ilo)*ldda + i__ - 1, ldda,
d_A, d_A_offset + N*ldda, ldda,
d_t, d_t_offset,
d_work, d_work_offset,
queue);
}
}
/* Use unblocked code to reduce the rest of the matrix */
if (!(nb < nbmin || nb >= nh))
magma_sgetmatrix( n, n-i__+1,
d_A, d_A_offset + (i__-ilo)*ldda, ldda,
a + (i__-1)*(lda), lda, queue );
lapackf77_sgehd2(&n, &i__, &ihi, a, &lda, &tau[1], work, &iinfo);
work[0] = MAGMA_S_MAKE( iws, 0 );
magma_free( da );
magma_free_cpu(t);
return *info;
} /* magma_sgehrd */
/**
Purpose
=======
CLAHEF computes a partial factorization of a complex Hermitian
matrix A using the Bunch-Kaufman diagonal pivoting method. The
partial factorization has the form:
A = ( I U12 ) ( A11 0 ) ( I 0 ) if UPLO = 'U', or:
( 0 U22 ) ( 0 D ) ( U12' U22' )
A = ( L11 0 ) ( D 0 ) ( L11' L21' ) if UPLO = 'L'
( L21 I ) ( 0 A22 ) ( 0 I )
where the order of D is at most NB. The actual order is returned in
the argument KB, and is either NB or NB-1, or N if N <= NB.
Note that U' denotes the conjugate transpose of U.
CLAHEF is an auxiliary routine called by CHETRF. It uses blocked code
(calling Level 3 BLAS) to update the submatrix A11 (if UPLO = 'U') or
A22 (if UPLO = 'L').
Arguments
---------
@param[in]
UPLO CHARACTER
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = 'U': Upper triangular
- = 'L': Lower triangular
@param[in]
N INTEGER
The order of the matrix A. N >= 0.
@param[in]
NB INTEGER
The maximum number of columns of the matrix A that should be
factored. NB should be at least 2 to allow for 2-by-2 pivot
blocks.
@param[out]
KB INTEGER
The number of columns of A that were actually factored.
KB is either NB-1 or NB, or N if N <= NB.
@param[in,out]
A COMPLEX 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, A contains details of the partial factorization.
@param[in]
LDA INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
ipiv INTEGER array, dimension (N)
Details of the interchanges and the block structure of D.
If UPLO = 'U', only the last KB elements of ipiv are set;
if UPLO = 'L', only the first KB elements are set.
\n
If ipiv(k) > 0, then rows and columns k and ipiv(k) were
interchanged and D(k,k) is a 1-by-1 diagonal block.
If UPLO = 'U' and ipiv(k) = ipiv(k-1) < 0, then rows and
columns k-1 and -ipiv(k) were interchanged and D(k-1:k,k-1:k)
is a 2-by-2 diagonal block. If UPLO = 'L' and ipiv(k) =
ipiv(k+1) < 0, then rows and columns k+1 and -ipiv(k) were
interchanged and D(k:k+1,k:k+1) is a 2-by-2 diagonal block.
@param[out]
W (workspace) COMPLEX array, dimension (LDW,NB)
@param[in]
LDW INTEGER
The leading dimension of the array W. LDW >= max(1,N).
@param[out]
INFO INTEGER
- = 0: successful exit
- > 0: if INFO = k, D(k,k) is exactly zero. The factorization
has been completed, but the block diagonal matrix D is
exactly singular.
@ingroup magma_chetrf_comp
********************************************************************/
extern "C" magma_int_t
magma_clahef_gpu(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb, magma_int_t *kb,
magmaFloatComplex *hA, magma_int_t lda,
magmaFloatComplex_ptr dA, size_t dA_offset, magma_int_t ldda,
magma_int_t *ipiv,
magmaFloatComplex_ptr dW, size_t dW_offset, magma_int_t lddw,
magma_queue_t queue,
magma_int_t *info)
{
/* .. Parameters .. */
float d_one = 1.0;
float d_zero = 0.0;
float d_eight = 8.0;
float d_seven = 7.0;
#if defined(PRECISION_c)
float f_zero = 0.0;
#endif
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_mone = -MAGMA_C_ONE;
magma_int_t upper = (uplo == MagmaUpper);
magma_int_t ione = 1;
/* .. Local Scalars .. */
magma_int_t imax = 0, jmax = 0, kk, kkW, kp, kstep, iinfo;
float abs_akk, alpha, colmax, R1, rowmax;
magmaFloatComplex Zimax, Z;
#define dA(i, j) dA, dA_offset + (j)*ldda + (i)
#define dW(i, j) dW, dW_offset + (j)*lddw + (i)
#define A(i, j) (hA + (j)*lda + (i))
/* .. Executable Statements .. */
*info = 0;
/* Initialize alpha for use in choosing pivot block size. */
alpha = ( d_one+sqrt( d_seven ) ) / d_eight;
magma_event_t event = NULL;
if( upper ) {
/* Factorize the trailing columns of A using the upper triangle
of A and working backwards, and compute the matrix W = U12*D
for use in updating A11 (note that conjg(W) is actually stored)
K is the main loop index, decreasing from N in steps of 1 or 2
KW is the column of W which corresponds to column K of A */
int k, kw = 0;
for (k = n-1; k+1 > max(n-nb+1, nb); k -= kstep) {
kw = nb - (n-k);
/* Copy column K of A to column KW of W and update it */
magma_ccopy( k+1, dA( 0, k ), 1, dW( 0, kw ), 1, queue );
// set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
if (k+1 < n) {
magma_cgemv( MagmaNoTrans, k+1, n-(k+1), c_mone, dA( 0, k+1 ), ldda,
dW( k, kw+1 ), lddw, c_one, dW( 0, kw ), ione, queue );
// set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k+ kw*lddw+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
}
kstep = 1;
/* Determine rows and columns to be interchanged and whether
a 1-by-1 or 2-by-2 pivot block will be used */
magma_cgetvector_async( 1, dW( k, kw ), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
abs_akk = fabs( MAGMA_C_REAL( Z ) );
/* imax is the row-index of the largest off-diagonal element in
column K, and colmax is its absolute value */
if( k > 0 ) {
// magma is one-base
imax = magma_icamax( k, dW( 0, kw ), 1, queue ) - 1;
magma_cgetvector( 1, dW( imax, kw ), 1, &Z, 1, queue );
colmax = MAGMA_C_ABS1( Z );
} else {
colmax = d_zero;
}
if( max( abs_akk, colmax ) == 0.0 ) {
/* Column K is zero: set INFO and continue */
if ( *info == 0 ) *info = k;
kp = k;
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(k+ k*ldda+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(k+ k*ldda+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
} else {
if( abs_akk >= alpha*colmax ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else {
/* Copy column imax to column KW-1 of W and update it */
magma_ccopy( imax+1, dA( 0, imax ), 1, dW( 0, kw-1 ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#endif
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( k-imax, dA(imax,imax+1), ldda, dW(imax+1,kw-1), 1, queue );
#else
magma_ccopy( k-imax, dA(imax,imax+1), ldda, dW(imax+1,kw-1), 1, queue );
#endif
if( k+1 < n ) {
magma_cgemv( MagmaNoTrans, k+1, n-(k+1), c_mone,
dA( 0, k+1 ), ldda, dW( imax, kw+1 ), lddw,
c_one, dW( 0, kw-1 ), ione, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(imax+ (kw-1)*lddw+dW_offset)+1, 1, queue, &event );
#endif
}
magma_cgetvector_async( 1, dW( imax, kw-1 ), 1, &Zimax, 1, queue, &event );
magma_queue_sync( queue );
/* jmax is the column-index of the largest off-diagonal
element in row imax, and rowmax is its absolute value */
jmax = imax + magma_icamax( k-imax, dW( imax+1, kw-1 ), 1, queue );
magma_cgetvector( 1, dW( jmax, kw-1 ), 1, &Z, 1, queue );
rowmax = MAGMA_C_ABS1( Z );
if ( imax > 0 ) {
// magma is one-base
jmax = magma_icamax( imax, dW( 0, kw-1 ), 1, queue ) - 1;
magma_cgetvector( 1, dW( jmax, kw-1 ), 1, &Z, 1, queue );
rowmax = max( rowmax, MAGMA_C_ABS1( Z ) );
}
if( abs_akk >= alpha*colmax*( colmax / rowmax ) ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else if ( fabs( MAGMA_C_REAL( Zimax ) ) >= alpha*rowmax ) {
/* interchange rows and columns K and imax, use 1-by-1
pivot block */
kp = imax;
/* copy column KW-1 of W to column KW */
magma_ccopy( k+1, dW( 0, kw-1 ), 1, dW( 0, kw ), 1, queue );
} else {
/* interchange rows and columns K-1 and imax, use 2-by-2
pivot block */
kp = imax;
kstep = 2;
}
}
kk = k - kstep + 1;
kkW = nb - (n - kk);
/* Updated column kp is already stored in column kkW of W */
if( kp != kk ) {
/* Interchange rows kk and kp in last kk columns of A and W */
// note: row-swap A(:,kk)
magmablas_cswap( n-kk, dA( kk, kk ), ldda, dA( kp, kk ), ldda, queue );
magmablas_cswap( n-kk, dW( kk, kkW), lddw, dW( kp, kkW), lddw, queue );
/* Copy non-updated column kk to column kp */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( kk-kp-1, dA( kp+1, kk ), 1, dA( kp, kp+1 ), ldda, queue );
#else
magma_ccopy( kk-kp-1, dA( kp+1, kk ), 1, dA( kp, kp+1 ), ldda, queue );
#endif
// now A(kp,kk) should be A(kk,kk), and copy to A(kp,kp)
magma_ccopy( kp+1, dA( 0, kk ), 1, dA( 0, kp ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(kp+ kp*ldda+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(kp+ kp*ldda+dA_offset)+1, 1, queue, &event );
#endif
}
if( kstep == 1 ) {
/* 1-by-1 pivot block D(k): column KW of W now holds
W(k) = U(k)*D(k)
where U(k) is the k-th column of U
Store U(k) in column k of A */
magma_ccopy( k+1, dW( 0, kw ), 1, dA( 0, k ), 1, queue );
if ( k > 0 ) {
magma_cgetvector_async( 1, dA( k, k ), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
R1 = d_one / MAGMA_C_REAL( Z );
magma_csscal( k, R1, dA( 0, k ), 1, queue );
/* Conjugate W(k) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( k, dW( 0, kw ), 1, dW( 0, kw ), 1, queue );
#endif
}
} else {
/* 2-by-2 pivot block D(k): columns KW and KW-1 of W now hold
( W(k-1) W(k) ) = ( U(k-1) U(k) )*D(k)
where U(k) and U(k-1) are the k-th and (k-1)-th columns of U */
if( k > 1 ) {
/* Store U(k) and U(k-1) in columns k and k-1 of A */
magmablas_clascl_2x2( MagmaUpper,
k-1, dW(0, kw-1), lddw, dA(0,k-1), ldda, &iinfo, queue );
}
/* Copy D(k) to A */
magma_ccopymatrix( 2, 2, dW( k-1, kw-1 ), lddw, dA( k-1, k-1 ), ldda, queue );
/* Conjugate W(k) and W(k-1) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( k, dW( 0, kw ), 1, dW( 0, kw ), 1, queue );
magmablas_clacpy_cnjg( k-1, dW( 0, kw-1 ), 1, dW( 0, kw-1 ), 1, queue );
#endif
}
}
/* Store details of the interchanges in ipiv */
if( kstep == 1 ) {
ipiv[ k ] = 1+kp;
} else {
ipiv[ k ] = -(1+kp);
ipiv[ k-1 ] = -(1+kp);
}
}
/* Update the upper triangle of A11 (= A(1:k,1:k)) as
A11 := A11 - U12*D*U12' = A11 - U12*W'
computing blocks of NB columns at a time (note that conjg(W) is
actually stored) */
kw = nb - (n-k);
for (int j = ( k / nb )*nb; j >= 0; j -= nb ) {
int jb = min( nb, k-j+1 );
#ifdef SYMMETRIC_UPDATE
/* Update the upper triangle of the diagonal block */
for (int jj = j; jj < j + jb; jj++) {
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#endif
magma_cgemv( MagmaNoTrans, jj-j+1, n-(k+1), c_mone,
dA( j, k+1 ), ldda, dW( jj, kw+1 ), lddw, c_one,
dA( j, jj ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(jj+ jj*ldda+dA_offset)+1, 1, queue, &event );
#endif
}
/* Update the rectangular superdiagonal block */
magma_cgemm( MagmaNoTrans, MagmaTrans, j, jb, n-(k+1),
c_mone, dA( 0, k+1 ), ldda, dW( j, kw+1 ), lddw,
c_one, dA( 0, j ), ldda, queue );
#else
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#endif
magma_cgemm( MagmaNoTrans, MagmaTrans, j+jb, jb, n-(k+1),
c_mone, dA( 0, k+1 ), ldda,
dW( j, kw+1 ), lddw,
c_one, dA( 0, j ), ldda, queue );
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j+ j*ldda+dA_offset)+1, 2*(1+ldda), queue );
#endif
#endif
}
/* Put U12 in standard form by partially undoing the interchanges in columns k+1:n */
for (int j = k+1; j < n;)
{
int jj = j;
int jp = ipiv[ j ];
if( jp < 0 ) {
jp = -jp;
j = j + 1;
}
j = j + 1;
jp = jp - 1;
if( jp != jj && j < n )
magmablas_cswap( n-j, dA( jp, j ), ldda, dA( jj, j ), ldda, queue );
}
// copying the panel back to CPU
magma_cgetmatrix_async( n, n-(k+1), dA(0,k+1), ldda, A(0,k+1), lda, queue, &event );
magma_queue_sync( queue );
/* Set KB to the number of columns factorized */
*kb = n - (k+1);
} else {
/* Factorize the leading columns of A using the lower triangle
of A and working forwards, and compute the matrix W = L21*D
for use in updating A22 (note that conjg(W) is actually stored)
K is the main loop index, increasing from 1 in steps of 1 or 2 */
int k;
for (k = 0; k < min(nb-1,n); k += kstep) {
/* Copy column K of A to column K of W and update it */
/* -------------------------------------------------------------- */
magma_ccopy( n-k, dA( k, k ), 1, dW( k, k ), 1, queue );
// set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
/* -------------------------------------------------------------- */
magma_cgemv( MagmaNoTrans, n-k, k, c_mone, dA( k, 0 ), ldda,
dW( k, 0 ), lddw, c_one, dW( k, k ), ione, queue );
// re-set imaginary part of diagonal to be zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*(k*lddw+k+dW_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
kstep = 1;
/* Determine rows and columns to be interchanged and whether
a 1-by-1 or 2-by-2 pivot block will be used */
magma_cgetvector_async( 1, dW( k, k ), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
abs_akk = fabs( MAGMA_C_REAL( Z ) );
/* imax is the row-index of the largest off-diagonal element in
column K, and colmax is its absolute value */
if( k < n-1 ) {
// magmablas is one-base
imax = k + magma_icamax( n-k-1, dW(k+1,k), 1, queue );
magma_cgetvector( 1, dW( imax,k ), 1, &Z, 1, queue );
colmax = MAGMA_C_ABS1( Z );
} else {
colmax = d_zero;
}
if ( max( abs_akk, colmax ) == 0.0 ) {
/* Column K is zero: set INFO and continue */
if( *info == 0 ) *info = k;
kp = k;
// make sure the imaginary part of diagonal is zero
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dA, 2*(k*ldda+k+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dA, 2*(k*ldda+k+dA_offset)+1, 1, queue, &event );
magma_queue_sync( queue );
#endif
} else {
if ( abs_akk >= alpha*colmax ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else {
/* Copy column imax to column K+1 of W and update it */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( imax-k, dA(imax,k), ldda, dW(k,k+1), 1, queue );
#else
magma_ccopy( imax-k, dA( imax, k ), ldda, dW( k, k+1 ), 1, queue );
#endif
magma_ccopy( n-imax, dA( imax, imax ), 1, dW( imax, k+1 ), 1, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
magma_cgemv( MagmaNoTrans, n-k, k, c_mone, dA( k, 0 ), ldda,
dW( imax, 0 ), lddw, c_one, dW( k, k+1 ), ione, queue );
#if defined(PRECISION_z)
magma_dsetvector_async( 1, &d_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#elif defined(PRECISION_c)
magma_ssetvector_async( 1, &f_zero, 1,
dW, 2*((k+1)*lddw+imax+dW_offset)+1, 1, queue, &event);
magma_queue_sync( queue );
#endif
magma_cgetvector_async( 1, dW(imax,k+1), 1, &Zimax, 1, queue, &event);
magma_queue_sync( queue );
/* jmax is the column-index of the largest off-diagonal
element in row imax, and rowmax is its absolute value */
// magmablas is one-base
jmax = k-1 + magma_icamax( imax-k, dW(k, k+1), 1, queue );
magma_cgetvector( 1, dW(jmax,k+1), 1, &Z, 1, queue );
rowmax = MAGMA_C_ABS1( Z );
if( imax < n-1 ) {
// magmablas is one-base
jmax = imax + magma_icamax( (n-1)-imax, dW(imax+1,k+1), 1, queue);
magma_cgetvector( 1, dW(jmax,k+1), 1, &Z, 1, queue );
rowmax = max( rowmax, MAGMA_C_ABS1( Z ) );
}
if( abs_akk >= alpha*colmax*( colmax / rowmax ) ) {
/* no interchange, use 1-by-1 pivot block */
kp = k;
} else if( fabs( MAGMA_C_REAL( Zimax ) ) >= alpha*rowmax ) {
/* interchange rows and columns K and imax, use 1-by-1
pivot block */
kp = imax;
/* copy column K+1 of W to column K */
magma_ccopy( n-k, dW( k, k+1 ), 1, dW( k, k ), 1, queue );
} else {
/* interchange rows and columns K+1 and imax, use 2-by-2
pivot block */
kp = imax;
kstep = 2;
}
}
kk = k + kstep - 1;
/* Updated column kp is already stored in column kk of W */
if( kp != kk ) {
/* Copy non-updated column kk to column kp */
/* ------------------------------------------------------------------ */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( kp-kk, dA( kk, kk ), 1, dA( kp, kk ), ldda, queue );
#else
magma_ccopy( kp-kk, dA( kk, kk ), 1, dA( kp, kk ), ldda, queue );
#endif
if ( kp < n ) {
magma_ccopy( n-kp, dA( kp, kk), 1, dA( kp, kp ), 1, queue );
}
/* ------------------------------------------------------------------ */
/* Interchange rows kk and kp in first kk columns of A and W */
magmablas_cswap( kk+1, dA( kk, 0 ), ldda, dA( kp, 0 ), ldda, queue );
magmablas_cswap( kk+1, dW( kk, 0 ), lddw, dW( kp, 0 ), lddw, queue );
}
if ( kstep == 1 ) {
/* 1-by-1 pivot block D(k): column k of W now holds
W(k) = L(k)*D(k)
where L(k) is the k-th column of L
Store L(k) in column k of A */
magma_ccopy( n-k, dW( k, k ), 1, dA( k, k ), 1, queue );
if ( k < n-1 ) {
magma_cgetvector_async( 1, dA(k,k), 1, &Z, 1, queue, &event );
magma_queue_sync( queue );
R1 = d_one / MAGMA_C_REAL( Z );
magma_csscal((n-1)-k, R1, dA( k+1,k ), 1, queue);
/* Conjugate W(k) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( (n-1)-k, dW( k+1,k ), 1, dW( k+1,k ), 1, queue );
#endif
}
} else {
/* 2-by-2 pivot block D(k): columns k and k+1 of W now hold
( W(k) W(k+1) ) = ( L(k) L(k+1) )*D(k)
where L(k) and L(k+1) are the k-th and (k+1)-th columns
of L */
magmablas_clascl_2x2( MagmaLower,
n-(k+2), dW(k,k), lddw, dA(k+2,k), ldda, &iinfo,
queue );
/* Copy D(k) to A */
magma_ccopymatrix( 2, 2, dW( k, k ), lddw, dA( k, k ), ldda, queue );
/* Conjugate W(k) and W(k+1) */
#if defined(PRECISION_z) || defined(PRECISION_c)
magmablas_clacpy_cnjg( (n-1)-k, dW( k+1,k ), 1, dW( k+1,k ), 1, queue );
magmablas_clacpy_cnjg( (n-1)-k-1, dW( k+2,k+1), 1, dW( k+2,k+1 ), 1, queue );
#endif
}
}
/* Store details of the interchanges in ipiv */
if ( kstep == 1 ) {
ipiv[k] = kp+1;
} else {
ipiv[k] = -kp-1;
ipiv[k+1] = -kp-1;
}
}
/* Update the lower triangle of A22 (= A(k:n,k:n)) as
A22 := A22 - L21*D*L21' = A22 - L21*W'
computing blocks of NB columns at a time (note that conjg(W) is
actually stored) */
for( int j = k; j < n; j += nb ) {
int jb = min( nb, n-j );
/* Update the lower triangle of the diagonal block */
#ifdef SYMMETRIC_UPDATE
for (int jj = j; jj < j + jb; jj++) {
int jnb = j + jb - jj;
/* -------------------------------------------------------- */
magma_cgemv( MagmaNoTrans, jnb, k, c_mone, dA( jj, 0 ), ldda,
dW( jj, 0 ), lddw, c_one, dA( jj, jj ), ione, queue );
/* -------------------------------------------------------- */
}
/* Update the rectangular subdiagonal block */
if( j+jb < n ) {
int nk = n - (j+jb);
/* -------------------------------------------- */
magma_cgemm( MagmaNoTrans, MagmaTrans, nk, jb, k,
c_mone, dA( j+jb, 0 ), ldda,
dW( j, 0 ), lddw,
c_one, dA( j+jb, j ), ldda, queue );
/* ------------------------------------------- */
}
#else
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#endif
magma_cgemm( MagmaNoTrans, MagmaTrans, n-j, jb, k,
c_mone, dA( j, 0 ), ldda,
dW( j, 0 ), lddw,
c_one, dA( j, j ), ldda, queue );
#if defined(PRECISION_z)
magmablas_dlaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#elif defined(PRECISION_c)
magmablas_slaset(MagmaUpperLower, 1, jb,
0, 0, dA, 2*(j*ldda+j+dA_offset)+1, 2*(1+ldda), queue );
#endif
#endif
}
/* Put L21 in standard form by partially undoing the interchanges
in columns 1:k-1 */
for (int j = k; j > 0;) {
int jj = j;
int jp = ipiv[j-1];
if( jp < 0 ) {
jp = -jp;
j--;
}
j--;
if ( jp != jj && j >= 1 ) {
magmablas_cswap( j, dA( jp-1,0 ), ldda, dA( jj-1,0 ), ldda, queue );
}
}
// copying the panel back to CPU
magma_cgetmatrix_async( n, k, dA(0,0), ldda, A(0,0), lda, queue, &event );
magma_queue_sync( queue );
/* Set KB to the number of columns factorized */
*kb = k;
}
return *info;
/* End of CLAHEF */
}
extern "C" magma_int_t
magma_sorgqr(magma_int_t m, magma_int_t n, magma_int_t k,
float *A, magma_int_t lda,
float *tau,
float *dT, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.1) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
December 2013
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.
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) REAL 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 SGEQRF_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) REAL array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by SGEQRF_GPU.
DT (input) REAL 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_sgeqrf_gpu.
NB (input) 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.
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)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_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;
float *dA, *dV, *dW;
float *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_smalloc( &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_smalloc_cpu( &work, lwork );
if (work == NULL) {
magma_free( dA );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
float *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 sorgqr is slow for
// k smaller than the sorgqr's blocking size (new version can be up to 60x faster)
lapackf77_sorgqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_slacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_slaset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_slarfb( 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_ssetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_slaset( 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_slaset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_ssetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
// set panel to identity
magmablas_slaset( MagmaUpperLower, i, ib, dA(0, i), ldda );
magmablas_slaset_identity( mi, ib, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_slarfb_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_sgetmatrix( 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_sorgqr */
/**
Purpose
-------
SGEGQR orthogonalizes the N vectors given by a real M-by-N matrix A:
A = Q * R.
On exit, if successful, the orthogonal vectors Q overwrite A
and R is given in work (on the CPU memory).
The routine is designed for tall-and-skinny matrices: M >> N, N <= 128.
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
Arguments
---------
@param[in]
ikind INTEGER
Several versions are implemented indiceted by the ikind value:
1: This version uses normal equations and SVD in an iterative process
that makes the computation numerically accurate.
2: This version uses a standard LAPACK-based orthogonalization through
MAGMA's QR panel factorization (magma_sgeqr2x3_gpu) and magma_sorgqr
3: MGS
4. Cholesky QR [ Note: this method uses the normal equations which
squares the condition number of A, therefore
||I - Q'Q|| < O(eps cond(A)^2) ]
@param[in]
m INTEGER
The number of rows of the matrix A. m >= n >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. 128 >= n >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (ldda,n)
On entry, the m-by-n matrix A.
On exit, the m-by-n matrix Q with orthogonal columns.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,m).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param
dwork (GPU workspace) REAL array, dimension:
n^2 for ikind = 1
3 n^2 + min(m, n) + 2 for ikind = 2
0 (not used) for ikind = 3
n^2 for ikind = 4
@param[out]
work (CPU workspace) REAL array, dimension 3 n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work should be in pinned memory.
@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.
@ingroup magma_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sgegqr_gpu( magma_int_t ikind, magma_int_t m, magma_int_t n,
float *dA, magma_int_t ldda,
float *dwork, float *work,
magma_int_t *info )
{
#define work(i_,j_) (work + (i_) + (j_)*n)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magma_int_t i = 0, j, k, n2 = n*n;
magma_int_t ione = 1;
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (ikind < 1 || ikind > 4) {
*info = -1;
} else if (m < 0 || m < n) {
*info = -2;
} else if (n < 0 || n > 128) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (ikind == 1) {
// === Iterative, based on SVD ============================================================
float *U, *VT, *vt, *R, *G, *hwork, *tau;
float *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_smalloc_cpu( &hwork, 32 + 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork=n*n+32; // First part f hwork; used as workspace in svd
U = hwork + n*n + 32; // Size n*n
S = (float *)(U+n*n); // Size n
tau = U + n*n + n; // Size n
#if defined(PRECISION_c) || defined(PRECISION_z)
float *rwork;
magma_smalloc_cpu( &rwork, 5*n);
if ( rwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
do {
i++;
magma_sgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_sgetmatrix(n, n, dwork, n, G, n);
#if defined(PRECISION_s) || defined(PRECISION_d)
lapackf77_sgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
#else
lapackf77_sgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, rwork, info);
#endif
mins = 100.f, maxs = 0.f;
for (k=0; k < n; k++) {
S[k] = magma_ssqrt( S[k] );
if (S[k] < mins) mins = S[k];
if (S[k] > maxs) maxs = S[k];
}
for (k=0; k < n; k++) {
vt = VT + k*n;
for (j=0; j < n; j++)
vt[j] *= S[j];
}
lapackf77_sgeqrf(&n, &n, VT, &n, tau, hwork, &lwork, info);
if (i == 1)
blasf77_scopy(&n2, VT, &ione, R, &ione);
else
blasf77_strmm("l", "u", "n", "n", &n, &n, &c_one, VT, &n, R, &n);
magma_ssetmatrix(n, n, VT, n, dwork, n);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
if (mins > 0.00001f)
cn = maxs/mins;
//fprintf(stderr, "Iteration %d, cond num = %f \n", i, cn);
} while (cn > 10.f);
magma_free_cpu( hwork );
#if defined(PRECISION_c) || defined(PRECISION_z)
magma_free_cpu( rwork );
#endif
// ================== end of ikind == 1 ===================================================
}
else if (ikind == 2) {
// ================== LAPACK based ===================================================
magma_int_t min_mn = min(m, n);
magma_int_t nb = n;
float *dtau = dwork + 2*n*n, *d_T = dwork, *ddA = dwork + n*n;
float *tau = work+n*n;
magmablas_slaset( MagmaFull, n, n, c_zero, c_zero, d_T, n );
magma_sgeqr2x3_gpu(m, n, dA, ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn+2*n*n), info);
magma_sgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_sgetmatrix( n, n, ddA, n, work, n);
magma_sorgqr_gpu( m, n, n, dA, ldda, tau, d_T, nb, info );
// ================== end of ikind == 2 ===================================================
}
else if (ikind == 3) {
// ================== MGS ===================================================
for(magma_int_t j = 0; j<n; j++){
for(magma_int_t i = 0; i<j; i++){
*work(i, j) = magma_sdot(m, dA(0,i), 1, dA(0,j), 1);
magma_saxpy(m, -(*work(i,j)), dA(0,i), 1, dA(0,j), 1);
}
for(magma_int_t i = j; i<n; i++)
*work(i, j) = MAGMA_S_ZERO;
//*work(j,j) = MAGMA_S_MAKE( magma_snrm2(m, dA(0,j), 1), 0. );
*work(j,j) = magma_sdot(m, dA(0,j), 1, dA(0,j), 1);
*work(j,j) = MAGMA_S_MAKE( sqrt(MAGMA_S_REAL( *work(j,j) )), 0.);
magma_sscal(m, 1./ *work(j,j), dA(0,j), 1);
}
// ================== end of ikind == 3 ===================================================
}
else if (ikind == 4) {
// ================== Cholesky QR ===================================================
magma_sgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_sgetmatrix(n, n, dwork, n, work, n);
lapackf77_spotrf("u", &n, work, &n, info);
magma_ssetmatrix(n, n, work, n, dwork, n);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
// ================== end of ikind == 4 ===================================================
}
return *info;
} /* magma_sgegqr_gpu */
/**
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.
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 REAL 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 SGEQRF_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 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 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_sgeqrf_gpu.
@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_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sorgqr(
magma_int_t m, magma_int_t n, magma_int_t k,
float *A, magma_int_t lda,
float *tau,
magmaFloat_ptr dT, magma_int_t nb,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_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;
float *dA, *dV, *dW;
float *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;
}
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
// 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_smalloc( &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_smalloc_cpu( &work, lwork );
if (work == NULL) {
magma_free( dA );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
float *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 sorgqr is slow for
// k smaller than the sorgqr's blocking size (new version can be up to 60x faster)
lapackf77_sorgqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_slacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_slaset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_slarfb( 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_ssetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_slaset( 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_slaset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_ssetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
// set panel to identity
magmablas_slaset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda );
magmablas_slaset( MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_slarfb_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_sgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda);
}
magma_queue_destroy( stream );
magma_free( dA );
magma_free_cpu( work );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_sorgqr */
extern "C" magma_int_t
magma_spidr(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix *x,
magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = MAGMA_NOTCONVERGED;
// prepare solver feedback
solver_par->solver = Magma_PIDR;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
solver_par->init_res = 0.0;
solver_par->final_res = 0.0;
solver_par->iter_res = 0.0;
solver_par->runtime = 0.0;
// constants
const float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_n_one = MAGMA_S_NEG_ONE;
// internal user parameters
const magma_int_t smoothing = 1; // 0 = disable, 1 = enable
const float angle = 0.7; // [0-1]
// local variables
magma_int_t iseed[4] = {0, 0, 0, 1};
magma_int_t dof;
magma_int_t s;
magma_int_t distr;
magma_int_t k, i, sk;
magma_int_t innerflag;
float residual;
float nrm;
float nrmb;
float nrmr;
float nrmt;
float rho;
float om;
float tt;
float tr;
float gamma;
float alpha;
float mkk;
float fk;
// matrices and vectors
magma_s_matrix dxs = {Magma_CSR};
magma_s_matrix dr = {Magma_CSR}, drs = {Magma_CSR};
magma_s_matrix dP = {Magma_CSR}, dP1 = {Magma_CSR};
magma_s_matrix dG = {Magma_CSR};
magma_s_matrix dU = {Magma_CSR};
magma_s_matrix dM = {Magma_CSR};
magma_s_matrix df = {Magma_CSR};
magma_s_matrix dt = {Magma_CSR};
magma_s_matrix dc = {Magma_CSR};
magma_s_matrix dv = {Magma_CSR};
magma_s_matrix dbeta = {Magma_CSR}, hbeta = {Magma_CSR};
magma_s_matrix dlu = {Magma_CSR};
// chronometry
real_Double_t tempo1, tempo2;
// initial s space
// TODO: add option for 's' (shadow space number)
// Hack: uses '--restart' option as the shadow space number.
// This is not a good idea because the default value of restart option is used to detect
// if the user provided a custom restart. This means that if the default restart value
// is changed then the code will think it was the user (unless the default value is
// also updated in the 'if' statement below.
s = 1;
if ( solver_par->restart != 50 ) {
if ( solver_par->restart > A.num_cols ) {
s = A.num_cols;
} else {
s = solver_par->restart;
}
}
solver_par->restart = s;
// set max iterations
solver_par->maxiter = min( 2 * A.num_cols, solver_par->maxiter );
// check if matrix A is square
if ( A.num_rows != A.num_cols ) {
//printf("Matrix A is not square.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
goto cleanup;
}
// |b|
nrmb = magma_snrm2( b.num_rows, b.dval, 1, queue );
if ( nrmb == 0.0 ) {
magma_sscal( x->num_rows, MAGMA_S_ZERO, x->dval, 1, queue );
info = MAGMA_SUCCESS;
goto cleanup;
}
// r = b - A x
CHECK( magma_svinit( &dr, Magma_DEV, b.num_rows, 1, c_zero, queue ));
CHECK( magma_sresidualvec( A, b, *x, &dr, &nrmr, queue ));
// |r|
solver_par->init_res = nrmr;
solver_par->final_res = solver_par->init_res;
solver_par->iter_res = solver_par->init_res;
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = (real_Double_t)nrmr;
}
// check if initial is guess good enough
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
goto cleanup;
}
// P = randn(n, s)
// P = ortho(P)
//---------------------------------------
// P = 0.0
CHECK( magma_svinit( &dP, Magma_CPU, A.num_cols, s, c_zero, queue ));
// P = randn(n, s)
distr = 3; // 1 = unif (0,1), 2 = unif (-1,1), 3 = normal (0,1)
dof = dP.num_rows * dP.num_cols;
lapackf77_slarnv( &distr, iseed, &dof, dP.val );
// transfer P to device
CHECK( magma_smtransfer( dP, &dP1, Magma_CPU, Magma_DEV, queue ));
magma_smfree( &dP, queue );
// P = ortho(P1)
if ( dP1.num_cols > 1 ) {
// P = magma_sqr(P1), QR factorization
CHECK( magma_sqr( dP1.num_rows, dP1.num_cols, dP1, dP1.ld, &dP, NULL, queue ));
} else {
// P = P1 / |P1|
nrm = magma_snrm2( dof, dP1.dval, 1, queue );
nrm = 1.0 / nrm;
magma_sscal( dof, nrm, dP1.dval, 1, queue );
CHECK( magma_smtransfer( dP1, &dP, Magma_DEV, Magma_DEV, queue ));
}
magma_smfree( &dP1, queue );
//---------------------------------------
// allocate memory for the scalar products
CHECK( magma_svinit( &hbeta, Magma_CPU, s, 1, c_zero, queue ));
CHECK( magma_svinit( &dbeta, Magma_DEV, s, 1, c_zero, queue ));
// smoothing enabled
if ( smoothing > 0 ) {
// set smoothing solution vector
CHECK( magma_smtransfer( *x, &dxs, Magma_DEV, Magma_DEV, queue ));
// set smoothing residual vector
CHECK( magma_smtransfer( dr, &drs, Magma_DEV, Magma_DEV, queue ));
}
// G(n,s) = 0
CHECK( magma_svinit( &dG, Magma_DEV, A.num_cols, s, c_zero, queue ));
// U(n,s) = 0
CHECK( magma_svinit( &dU, Magma_DEV, A.num_cols, s, c_zero, queue ));
// M(s,s) = I
CHECK( magma_svinit( &dM, Magma_DEV, s, s, c_zero, queue ));
magmablas_slaset( MagmaFull, s, s, c_zero, c_one, dM.dval, s, queue );
// f = 0
CHECK( magma_svinit( &df, Magma_DEV, dP.num_cols, 1, c_zero, queue ));
// t = 0
CHECK( magma_svinit( &dt, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// c = 0
CHECK( magma_svinit( &dc, Magma_DEV, dM.num_cols, 1, c_zero, queue ));
// v = 0
CHECK( magma_svinit( &dv, Magma_DEV, dr.num_rows, 1, c_zero, queue ));
// lu = 0
CHECK( magma_svinit( &dlu, Magma_DEV, A.num_rows, 1, c_zero, queue ));
//--------------START TIME---------------
// chronometry
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->timing[0] = 0.0;
}
om = MAGMA_S_ONE;
innerflag = 0;
// start iteration
do
{
solver_par->numiter++;
// new RHS for small systems
// f = P' r
magmablas_sgemv( MagmaConjTrans, dP.num_rows, dP.num_cols, c_one, dP.dval, dP.ld, dr.dval, 1, c_zero, df.dval, 1, queue );
// shadow space loop
for ( k = 0; k < s; ++k ) {
sk = s - k;
// f(k:s) = M(k:s,k:s) c(k:s)
magma_scopyvector( sk, &df.dval[k], 1, &dc.dval[k], 1, queue );
magma_strsv( MagmaLower, MagmaNoTrans, MagmaNonUnit, sk, &dM.dval[k*dM.ld+k], dM.ld, &dc.dval[k], 1, queue );
// v = r - G(:,k:s) c(k:s)
magma_scopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
magmablas_sgemv( MagmaNoTrans, dG.num_rows, sk, c_n_one, &dG.dval[k*dG.ld], dG.ld, &dc.dval[k], 1, c_one, dv.dval, 1, queue );
// preconditioning operation
// v = L \ v;
// v = U \ v;
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queue ));
// U(:,k) = om * v + U(:,k:s) c(k:s)
magmablas_sgemv( MagmaNoTrans, dU.num_rows, sk, c_one, &dU.dval[k*dU.ld], dU.ld, &dc.dval[k], 1, om, dv.dval, 1, queue );
magma_scopyvector( dU.num_rows, dv.dval, 1, &dU.dval[k*dU.ld], 1, queue );
// G(:,k) = A U(:,k)
CHECK( magma_s_spmv( c_one, A, dv, c_zero, dv, queue ));
solver_par->spmv_count++;
magma_scopyvector( dG.num_rows, dv.dval, 1, &dG.dval[k*dG.ld], 1, queue );
// bi-orthogonalize the new basis vectors
for ( i = 0; i < k; ++i ) {
// alpha = P(:,i)' G(:,k)
alpha = magma_sdot( dP.num_rows, &dP.dval[i*dP.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// alpha = alpha / M(i,i)
magma_sgetvector( 1, &dM.dval[i*dM.ld+i], 1, &mkk, 1, queue );
alpha = alpha / mkk;
// G(:,k) = G(:,k) - alpha * G(:,i)
magma_saxpy( dG.num_rows, -alpha, &dG.dval[i*dG.ld], 1, &dG.dval[k*dG.ld], 1, queue );
// U(:,k) = U(:,k) - alpha * U(:,i)
magma_saxpy( dU.num_rows, -alpha, &dU.dval[i*dU.ld], 1, &dU.dval[k*dU.ld], 1, queue );
}
// new column of M = P'G, first k-1 entries are zero
// M(k:s,k) = P(:,k:s)' G(:,k)
magmablas_sgemv( MagmaConjTrans, dP.num_rows, sk, c_one, &dP.dval[k*dP.ld], dP.ld, &dG.dval[k*dG.ld], 1, c_zero, &dM.dval[k*dM.ld+k], 1, queue );
// check M(k,k) == 0
magma_sgetvector( 1, &dM.dval[k*dM.ld+k], 1, &mkk, 1, queue );
if ( MAGMA_S_EQUAL(mkk, MAGMA_S_ZERO) ) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// beta = f(k) / M(k,k)
magma_sgetvector( 1, &df.dval[k], 1, &fk, 1, queue );
hbeta.val[k] = fk / mkk;
// check for nan
if ( magma_s_isnan( hbeta.val[k] ) || magma_s_isinf( hbeta.val[k] )) {
innerflag = 1;
info = MAGMA_DIVERGENCE;
break;
}
// r = r - beta * G(:,k)
magma_saxpy( dr.num_rows, -hbeta.val[k], &dG.dval[k*dG.ld], 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// |r|
nrmr = magma_snrm2( dr.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// x = x + beta * U(:,k)
magma_saxpy( x->num_rows, hbeta.val[k], &dU.dval[k*dU.ld], 1, x->dval, 1, queue );
// smoothing operation
//---------------------------------------
// t = rs - r
magma_scopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_sdot( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_sdot( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (t' * t)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_saxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_scopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_saxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_snrm2( drs.num_rows, drs.dval, 1, queue );
//---------------------------------------
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter) % solver_par->verbose == 0 ) {
solver_par->res_vec[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
s = k + 1; // for the x-update outside the loop
innerflag = 2;
info = MAGMA_SUCCESS;
break;
}
// non-last s iteration
if ( (k + 1) < s ) {
// f(k+1:s) = f(k+1:s) - beta * M(k+1:s,k)
magma_saxpy( sk-1, -hbeta.val[k], &dM.dval[k*dM.ld+(k+1)], 1, &df.dval[k+1], 1, queue );
}
}
// smoothing disabled
if ( smoothing <= 0 && innerflag != 1 ) {
// update solution approximation x
// x = x + U(:,1:s) * beta(1:s)
magma_ssetvector( s, hbeta.val, 1, dbeta.dval, 1, queue );
magmablas_sgemv( MagmaNoTrans, dU.num_rows, s, c_one, dU.dval, dU.ld, dbeta.dval, 1, c_one, x->dval, 1, queue );
}
// check convergence or iteration limit or invalid result of inner loop
if ( innerflag > 0 ) {
break;
}
// v = r
magma_scopyvector( dr.num_rows, dr.dval, 1, dv.dval, 1, queue );
// preconditioning operation
// v = L \ v;
// v = U \ v;
CHECK( magma_s_applyprecond_left( MagmaNoTrans, A, dv, &dlu, precond_par, queue ));
CHECK( magma_s_applyprecond_right( MagmaNoTrans, A, dlu, &dv, precond_par, queue ));
// t = A v
CHECK( magma_s_spmv( c_one, A, dv, c_zero, dt, queue ));
solver_par->spmv_count++;
// computation of a new omega
//---------------------------------------
// |t|
nrmt = magma_snrm2( dt.num_rows, dt.dval, 1, queue );
// t'r
tr = magma_sdot( dt.num_rows, dt.dval, 1, dr.dval, 1, queue );
// rho = abs(t' * r) / (|t| * |r|))
rho = MAGMA_D_ABS( MAGMA_S_REAL(tr) / (nrmt * nrmr) );
// om = (t' * r) / (|t| * |t|)
om = tr / (nrmt * nrmt);
if ( rho < angle ) {
om = (om * angle) / rho;
}
//---------------------------------------
if ( MAGMA_S_EQUAL(om, MAGMA_S_ZERO) ) {
info = MAGMA_DIVERGENCE;
break;
}
// update approximation vector
// x = x + om * v
magma_saxpy( x->num_rows, om, dv.dval, 1, x->dval, 1, queue );
// update residual vector
// r = r - om * t
magma_saxpy( dr.num_rows, -om, dt.dval, 1, dr.dval, 1, queue );
// smoothing disabled
if ( smoothing <= 0 ) {
// residual norm
nrmr = magma_snrm2( b.num_rows, dr.dval, 1, queue );
// smoothing enabled
} else {
// smoothing operation
//---------------------------------------
// t = rs - r
magma_scopyvector( drs.num_rows, drs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, dr.dval, 1, dt.dval, 1, queue );
// t't
// t'rs
tt = magma_sdot( dt.num_rows, dt.dval, 1, dt.dval, 1, queue );
tr = magma_sdot( dt.num_rows, dt.dval, 1, drs.dval, 1, queue );
// gamma = (t' * rs) / (|t| * |t|)
gamma = tr / tt;
// rs = rs - gamma * (rs - r)
magma_saxpy( drs.num_rows, -gamma, dt.dval, 1, drs.dval, 1, queue );
// xs = xs - gamma * (xs - x)
magma_scopyvector( dxs.num_rows, dxs.dval, 1, dt.dval, 1, queue );
magma_saxpy( dt.num_rows, c_n_one, x->dval, 1, dt.dval, 1, queue );
magma_saxpy( dxs.num_rows, -gamma, dt.dval, 1, dxs.dval, 1, queue );
// |rs|
nrmr = magma_snrm2( b.num_rows, drs.dval, 1, queue );
//---------------------------------------
}
// store current timing and residual
if ( solver_par->verbose > 0 ) {
tempo2 = magma_sync_wtime( queue );
if ( (solver_par->numiter) % solver_par->verbose == 0 ) {
solver_par->res_vec[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)nrmr;
solver_par->timing[(solver_par->numiter) / solver_par->verbose]
= (real_Double_t)tempo2 - tempo1;
}
}
// check convergence
if ( nrmr <= solver_par->atol ||
nrmr/nrmb <= solver_par->rtol ) {
info = MAGMA_SUCCESS;
break;
}
}
while ( solver_par->numiter + 1 <= solver_par->maxiter );
// smoothing enabled
if ( smoothing > 0 ) {
// x = xs
magma_scopyvector( x->num_rows, dxs.dval, 1, x->dval, 1, queue );
// r = rs
magma_scopyvector( dr.num_rows, drs.dval, 1, dr.dval, 1, queue );
}
// get last iteration timing
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t)tempo2 - tempo1;
//--------------STOP TIME----------------
// get final stats
solver_par->iter_res = nrmr;
CHECK( magma_sresidualvec( A, b, *x, &dr, &residual, queue ));
solver_par->final_res = residual;
// set solver conclusion
if ( info != MAGMA_SUCCESS && info != MAGMA_DIVERGENCE ) {
if ( solver_par->init_res > solver_par->final_res ) {
info = MAGMA_SLOW_CONVERGENCE;
}
}
cleanup:
// free resources
// smoothing enabled
if ( smoothing > 0 ) {
magma_smfree( &dxs, queue );
magma_smfree( &drs, queue );
}
magma_smfree( &dr, queue );
magma_smfree( &dP, queue );
magma_smfree( &dP1, queue );
magma_smfree( &dG, queue );
magma_smfree( &dU, queue );
magma_smfree( &dM, queue );
magma_smfree( &df, queue );
magma_smfree( &dt, queue );
magma_smfree( &dc, queue );
magma_smfree( &dv, queue );
magma_smfree(&dlu, queue);
magma_smfree( &dbeta, queue );
magma_smfree( &hbeta, queue );
solver_par->info = info;
return info;
/* magma_spidr */
}
int main(int argc, char **argv)
{
TESTING_INIT();
const float c_neg_one = MAGMA_S_NEG_ONE;
const magma_int_t ione = 1;
real_Double_t atomics_perf, atomics_time;
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, atomics_error, cublas_error, work[1];
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t N, lda, ldda, sizeA, sizeX, sizeY, blocks, ldwork;
magma_int_t incx = 1;
magma_int_t incy = 1;
magma_int_t nb = 64;
float alpha = MAGMA_S_MAKE( 1.5, -2.3 );
float beta = MAGMA_S_MAKE( -0.6, 0.8 );
float *A, *X, *Y, *Yatomics, *Ycublas, *Ymagma;
magmaFloat_ptr dA, dX, dY, dwork;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N MAGMA Gflop/s (ms) Atomics Gflop/s CUBLAS Gflop/s CPU Gflop/s MAGMA error Atomics CUBLAS\n");
printf("======================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = ((N + 31)/32)*32;
sizeA = N*lda;
sizeX = N*incx;
sizeY = N*incy;
gflops = FLOPS_SSYMV( N ) / 1e9;
TESTING_MALLOC_CPU( A, float, sizeA );
TESTING_MALLOC_CPU( X, float, sizeX );
TESTING_MALLOC_CPU( Y, float, sizeY );
TESTING_MALLOC_CPU( Yatomics, float, sizeY );
TESTING_MALLOC_CPU( Ycublas, float, sizeY );
TESTING_MALLOC_CPU( Ymagma, float, sizeY );
TESTING_MALLOC_DEV( dA, float, ldda*N );
TESTING_MALLOC_DEV( dX, float, sizeX );
TESTING_MALLOC_DEV( dY, float, sizeY );
blocks = (N + nb - 1) / nb;
ldwork = ldda*blocks;
TESTING_MALLOC_DEV( dwork, float, ldwork );
magmablas_slaset( MagmaFull, ldwork, 1, MAGMA_S_NAN, MAGMA_S_NAN, dwork, ldwork );
magmablas_slaset( MagmaFull, ldda, N, MAGMA_S_NAN, MAGMA_S_NAN, dA, ldda );
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &sizeA, A );
magma_smake_symmetric( N, A, lda );
// should not use data from the opposite triangle -- fill with NAN to check
magma_int_t N1 = N-1;
if ( opts.uplo == MagmaUpper ) {
lapackf77_slaset( "Lower", &N1, &N1, &MAGMA_S_NAN, &MAGMA_S_NAN, &A[1], &lda );
}
else {
lapackf77_slaset( "Upper", &N1, &N1, &MAGMA_S_NAN, &MAGMA_S_NAN, &A[lda], &lda );
}
lapackf77_slarnv( &ione, ISEED, &sizeX, X );
lapackf77_slarnv( &ione, ISEED, &sizeY, Y );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_ssetmatrix( N, N, A, lda, dA, ldda );
magma_ssetvector( N, X, incx, dX, incx );
magma_ssetvector( N, Y, incy, dY, incy );
cublas_time = magma_sync_wtime( 0 );
cublasSsymv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
cublas_time = magma_sync_wtime( 0 ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_sgetvector( N, dY, incy, Ycublas, incy );
/* =====================================================================
Performs operation using CUBLAS - using atomics
=================================================================== */
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_ALLOWED );
magma_ssetvector( N, Y, incy, dY, incy );
atomics_time = magma_sync_wtime( 0 );
cublasSsymv( opts.handle, cublas_uplo_const(opts.uplo),
N, &alpha, dA, ldda, dX, incx, &beta, dY, incy );
atomics_time = magma_sync_wtime( 0 ) - atomics_time;
atomics_perf = gflops / atomics_time;
magma_sgetvector( N, dY, incy, Yatomics, incy );
cublasSetAtomicsMode( opts.handle, CUBLAS_ATOMICS_NOT_ALLOWED );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_ssetvector( N, Y, incy, dY, incy );
magma_time = magma_sync_wtime( 0 );
if ( opts.version == 1 ) {
magmablas_ssymv_work( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy, dwork, ldwork, opts.queue );
}
else {
// non-work interface (has added overhead)
magmablas_ssymv( opts.uplo, N, alpha, dA, ldda, dX, incx, beta, dY, incy );
}
magma_time = magma_sync_wtime( 0 ) - magma_time;
magma_perf = gflops / magma_time;
magma_sgetvector( N, dY, incy, Ymagma, incy );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_ssymv( lapack_uplo_const(opts.uplo), &N, &alpha, A, &lda, X, &incx, &beta, Y, &incy );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Ymagma, &incy );
magma_error = lapackf77_slange( "M", &N, &ione, Ymagma, &N, work ) / N;
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Ycublas, &incy );
cublas_error = lapackf77_slange( "M", &N, &ione, Ycublas, &N, work ) / N;
blasf77_saxpy( &N, &c_neg_one, Y, &incy, Yatomics, &incy );
atomics_error = lapackf77_slange( "M", &N, &ione, Yatomics, &N, work ) / N;
bool ok = (magma_error < tol && cublas_error < tol && atomics_error < tol);
status += ! ok;
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e %s\n",
(int) N,
magma_perf, 1000.*magma_time,
atomics_perf, 1000.*atomics_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error, atomics_error,
(ok ? "ok" : "failed"));
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( X );
TESTING_FREE_CPU( Y );
TESTING_FREE_CPU( Ycublas );
TESTING_FREE_CPU( Yatomics );
TESTING_FREE_CPU( Ymagma );
TESTING_FREE_DEV( dA );
TESTING_FREE_DEV( dX );
TESTING_FREE_DEV( dY );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
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, 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 (workspace) REAL work space array on the GPU,
dimension (2*MIN(M, N) + (N+31)/32*32 )*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_sgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_sorgqr_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(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, lpanel;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork;
float *dV, *dW;
float *work, *panel;
*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;
}
// 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 CPU work space
// n*nb for sorgqr workspace
// (m - kk)*(n - kk) for last block's panel
lwork = n*nb;
lpanel = (m - kk)*(n - kk);
magma_smalloc_cpu( &work, lwork + lpanel );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
panel = work + lwork;
// Allocate work space on GPU
if (MAGMA_SUCCESS != magma_smalloc( &dV, ldda*nb )) {
magma_free_cpu( work );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// dT workspace has:
// 2*min(m,n)*nb for T and R^{-1} matrices from geqrf
// ((n+31)/32*32 )*nb for dW larfb workspace.
lddwork = min(m,n);
dW = dT + 2*lddwork*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;
magma_sgetmatrix( m_kk, k_kk,
dA(kk, kk), ldda, panel, m_kk );
lapackf77_sorgqr( &m_kk, &n_kk, &k_kk,
panel, &m_kk,
&tau[kk], work, &lwork, &iinfo );
magma_ssetmatrix( m_kk, n_kk,
panel, m_kk, dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_slaset( MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda );
}
if (kk > 0) {
// Use blocked code
// stream: copy Aii to V --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min( nb, k-i );
mi = m - i;
// Copy current panel on the GPU from dA to dV
magma_scopymatrix_async( mi, ib,
dA(i,i), ldda,
dV, ldda, stream );
// set panel to identity
magmablas_slaset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda );
magmablas_slaset( MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
}
magma_queue_sync( stream );
magmablasSetKernelStream( NULL );
magma_free( dV );
magma_free_cpu( work );
magma_queue_destroy( stream );
return *info;
} /* magma_sorgqr_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const float d_neg_one = MAGMA_D_NEG_ONE;
const float d_one = MAGMA_D_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
const float c_one = MAGMA_S_ONE;
const float c_zero = MAGMA_S_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float Anorm, error=0, error2=0;
float *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloat_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
opts.parse_opts( argc, argv );
magma_int_t status = 0;
float tol = opts.tolerance * lapackf77_slamch("E");
// version 3 can do either check
if (opts.check == 1 && opts.version == 1) {
opts.check = 2;
printf( "%% version 1 requires check 2 (solve A*x=b)\n" );
}
if (opts.check == 2 && opts.version == 2) {
opts.check = 1;
printf( "%% version 2 requires check 1 (R - Q^H*A)\n" );
}
printf( "%% version %d\n", (int) opts.version );
if ( opts.check == 1 ) {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |R - Q^H*A| |I - Q^H*Q|\n");
printf("%%==============================================================================\n");
}
else {
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |b - A*x|\n");
printf("%%===============================================================\n");
}
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min( M, N );
lda = M;
n2 = lda*N;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
nb = magma_get_sgeqrf_nb( M, N );
gflops = FLOPS_SGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_sgeqrf( &M, &N, NULL, &M, NULL, tmp, &lwork, &info );
lwork = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, float, min_mn );
TESTING_MALLOC_CPU( h_A, float, n2 );
TESTING_MALLOC_CPU( h_work, float, lwork );
TESTING_MALLOC_PIN( h_R, float, n2 );
TESTING_MALLOC_DEV( d_A, float, ldda*N );
if ( opts.version == 1 || opts.version == 3 ) {
size = (2*min(M, N) + magma_roundup( N, 32 ) )*nb;
TESTING_MALLOC_DEV( dT, float, size );
magmablas_slaset( MagmaFull, size, 1, c_zero, c_zero, dT, size );
}
/* Initialize the matrix */
lapackf77_slarnv( &ione, ISEED, &n2, h_A );
lapackf77_slacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
magma_ssetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
nb = magma_get_sgeqrf_nb( M, N );
gpu_time = magma_wtime();
if ( opts.version == 1 ) {
// stores dT, V blocks have zeros, R blocks inverted & stored in dT
magma_sgeqrf_gpu( M, N, d_A, ldda, tau, dT, &info );
}
else if ( opts.version == 2 ) {
// LAPACK complaint arguments
magma_sgeqrf2_gpu( M, N, d_A, ldda, tau, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// stores dT, V blocks have zeros, R blocks stored in dT
magma_sgeqrf3_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
return -1;
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
if ( opts.check == 1 && (opts.version == 2 || opts.version == 3) ) {
if ( opts.version == 3 ) {
// copy diagonal blocks of R back to A
for( int i=0; i < min_mn-nb; i += nb ) {
magma_int_t ib = min( min_mn-i, nb );
magmablas_slacpy( MagmaUpper, ib, ib, &dT[min_mn*nb + i*nb], nb, &d_A[ i + i*ldda ], ldda );
}
}
/* =====================================================================
Check the result, following zqrt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
Only for version 2, which has LAPACK complaint output.
Or for version 3, after restoring diagonal blocks of A above.
=================================================================== */
magma_sgetmatrix( M, N, d_A, ldda, h_R, lda );
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
float *Q, *R;
float *work;
TESTING_MALLOC_CPU( Q, float, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, float, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, float, min_mn );
// generate M by K matrix Q, where K = min(M,N)
lapackf77_slacpy( "Lower", &M, &min_mn, h_R, &lda, Q, &ldq );
lapackf77_sorgqr( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy K by N matrix R
lapackf77_slaset( "Lower", &min_mn, &N, &c_zero, &c_zero, R, &ldr );
lapackf77_slacpy( "Upper", &min_mn, &N, h_R, &lda, R, &ldr );
// error = || R - Q^H*A || / (N * ||A||)
blasf77_sgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, R, &ldr );
Anorm = lapackf77_slange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_slange( "1", &min_mn, &N, R, &ldr, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set R = I (K by K identity), then R = I - Q^H*Q
// error = || I - Q^H*Q || / N
lapackf77_slaset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, R, &ldr );
blasf77_ssyrk( "Upper", "Conj", &min_mn, &M, &d_neg_one, Q, &ldq, &d_one, R, &ldr );
error2 = safe_lapackf77_slansy( "1", "Upper", &min_mn, R, &ldr, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
else if ( opts.check == 2 && M >= N && (opts.version == 1 || opts.version == 3) ) {
/* =====================================================================
Check the result by solving consistent linear system, A*x = b.
Only for versions 1 & 3 with M >= N.
=================================================================== */
magma_int_t lwork2;
float *x, *b, *hwork;
magmaFloat_ptr d_B;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, float, N );
TESTING_MALLOC_CPU( b, float, M );
lapackf77_slarnv( &ione, ISEED, &N, x );
blasf77_sgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, float, M );
magma_ssetvector( M, b, 1, d_B, 1 );
if ( opts.version == 1 ) {
// allocate hwork
magma_sgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork2 = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lwork2 );
// solve linear system
magma_sgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork2, &info );
if (info != 0) {
printf("magma_sgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_sgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork2 = (magma_int_t)MAGMA_S_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, float, lwork2 );
// solve linear system
magma_sgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork2, &info );
if (info != 0) {
printf("magma_sgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
return -1;
}
magma_sgetvector( N, d_B, 1, x, 1 );
// compute r = Ax - b, saved in b
blasf77_sgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (max(m,n)*|A|*|x|)
float norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_slange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_slange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_slange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (max(M,N) * norm_A * norm_x);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_sgeqrf( &M, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_sgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check == 1 ) {
bool okay = (error < tol && error2 < tol);
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else if ( opts.check == 2 ) {
if ( M >= N ) {
bool okay = (error < tol);
status += ! okay;
printf( "%10.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf( "(error check only for M >= N)\n" );
}
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version == 1 || opts.version == 3 ) {
TESTING_FREE_DEV( dT );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
SSYTRD2_GPU reduces a real symmetric matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version passes a workspace that is used in an optimized
GPU matrix-vector product.
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]
dA REAL array on the GPU, dimension (LDDA,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]
ldda INTEGER
The leading dimension of the array A. LDDA >= 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]
A (workspace) REAL array, dimension (LDA,N)
On exit the diagonal, the upper part (if uplo=MagmaUpper)
or the lower part (if uplo=MagmaLower) are copies of DA
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@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 >= N*NB, where NB is the
optimal blocksize given by magma_get_ssytrd_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]
dwork (workspace) REAL array on the GPU, dim (MAX(1,LDWORK))
@param[in]
ldwork INTEGER
The dimension of the array DWORK.
LDWORK >= ldda*ceil(n/64) + 2*ldda*nb, where nb = magma_get_ssytrd_nb(n),
and 64 is for the blocksize of magmablas_ssymv.
@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_ssytrd2_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaFloat_ptr dA, magma_int_t ldda,
float *d, float *e, float *tau,
float *A, magma_int_t lda,
float *work, magma_int_t lwork,
magmaFloat_ptr dwork, magma_int_t ldwork,
magma_int_t *info)
{
#define A(i_, j_) ( A + (i_) + (j_)*lda )
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
/* Constants */
const float c_zero = MAGMA_S_ZERO;
const float c_neg_one = MAGMA_S_NEG_ONE;
const float c_one = MAGMA_S_ONE;
const float d_one = MAGMA_D_ONE;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t nb = magma_get_ssytrd_nb( n );
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;
bool upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -9;
} else if (lwork < nb*n && ! lquery) {
*info = -11;
} else if (ldwork < ldda*magma_ceildiv(n,64) + 2*ldda*nb) {
*info = -13;
}
/* Determine the block size. */
ldw = n;
lddw = ldda; // hopefully ldda is rounded up to multiple of 32; ldwork is in terms of ldda, so lddw can't be > ldda.
lwkopt = n * nb;
if (*info == 0) {
work[0] = magma_smake_lwork( 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;
}
// nx <= n is required
// use LAPACK for n < 3000, otherwise switch at 512
if (n < 3000)
nx = n;
else
nx = 512;
float *work2;
if (MAGMA_SUCCESS != magma_smalloc_cpu( &work2, n )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t queue = NULL;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
// clear out dwork in case it has NANs (used as y in ssymv)
// rest of dwork (used as work in magmablas_ssymv) doesn't need to be cleared
magmablas_slaset( MagmaFull, n, nb, c_zero, c_zero, dwork, lddw, queue );
if (upper) {
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
kk = n - magma_roundup( n - nx, 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 */
magma_sgetmatrix( i+nb, nb, dA(0, i), ldda, A(0, i), lda, queue );
magma_slatrd2( uplo, i+nb, nb, A(0, 0), lda, e, tau,
work, ldw, work2, n, dA(0, 0), ldda, dwork, lddw,
dwork + 2*lddw*nb, ldwork - 2*lddw*nb, 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, ldw, dwork, lddw, queue );
magma_ssyr2k( uplo, MagmaNoTrans, i, nb, c_neg_one,
dA(0, i), ldda, dwork, lddw,
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) {
*A(j-1,j) = MAGMA_S_MAKE( e[j - 1], 0 );
d[j] = MAGMA_S_REAL( *A(j, j) );
}
}
magma_sgetmatrix( kk, kk, dA(0, 0), ldda, A(0, 0), lda, queue );
/* Use CPU code to reduce the last or only block */
lapackf77_ssytrd( uplo_, &kk, A(0, 0), &lda, d, e, tau, work, &lwork, &iinfo );
magma_ssetmatrix( kk, kk, A(0, 0), lda, dA(0, 0), ldda, queue );
}
else {
/* 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 */
magma_sgetmatrix( n-i, nb, dA(i, i), ldda, A(i, i), lda, queue );
magma_slatrd2( uplo, n-i, nb, A(i, i), lda, &e[i], &tau[i],
work, ldw, work2, n, dA(i, i), ldda, dwork, lddw,
dwork + 2*lddw*nb, ldwork - 2*lddw*nb, queue );
/* 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, ldw, dwork, lddw, queue );
// cublas 6.5 crashes here if lddw % 32 != 0, e.g., N=250.
magma_ssyr2k( 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, queue );
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+nb; ++j) {
*A(j+1,j) = MAGMA_S_MAKE( e[j], 0 );
d[j] = MAGMA_S_REAL( *A(j, j) );
}
}
/* Use CPU code to reduce the last or only block */
magma_sgetmatrix( n-i, n-i, dA(i, i), ldda, A(i, i), 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_ssetmatrix( n-i, n-i, A(i, i), lda, dA(i, i), ldda, queue );
}
magma_free_cpu( work2 );
magma_queue_destroy( queue );
work[0] = magma_smake_lwork( lwkopt );
return *info;
} /* magma_ssytrd2_gpu */
