extern "C" magma_int_t
magma_c_applyprecond_right(
magma_c_sparse_matrix A,
magma_c_vector b,
magma_c_vector *x,
magma_c_preconditioner *precond,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
if ( precond->solver == Magma_JACOBI ) {
magma_ccopy( b.num_rows*b.num_cols, b.dval, 1, x->dval, 1 ); // x = b
}
else if ( precond->solver == Magma_ILU ||
( precond->solver == Magma_AILU && precond->maxiter == -1)) {
magma_capplycumilu_r( b, x, precond, queue );
}
else if ( precond->solver == Magma_ICC ||
( precond->solver == Magma_AICC && precond->maxiter == -1) ) {
magma_capplycumicc_r( b, x, precond, queue );
}
else if ( precond->solver == Magma_NONE ) {
magma_ccopy( b.num_rows*b.num_cols, b.dval, 1, x->dval, 1 ); // x = b
}
else {
printf( "error: preconditioner type not yet supported.\n" );
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_zjacobiiter_precond(
magma_z_sparse_matrix M, magma_z_vector *x,
magma_z_solver_par *solver_par, magma_z_preconditioner *precond,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// local variables
magmaDoubleComplex c_zero = MAGMA_Z_ZERO, c_one = MAGMA_Z_ONE,
c_mone = MAGMA_Z_NEG_ONE;
magma_int_t dofs = M.num_rows;
magma_int_t num_vecs = x->num_rows / dofs;
magma_z_vector swap;
for( magma_int_t i=0; i<solver_par->maxiter; i++ ) {
magma_z_spmv( c_mone, M, *x, c_zero, precond->work2, queue ); // t = - M * x
magma_zaxpy( num_vecs*dofs, c_one ,
precond->work1.dval, 1 , precond->work2.dval, 1 ); // t = t + c
// swap so that x again contains solution, and y is ready to be used
swap = *x;
*x = precond->work2;
precond->work2 = swap;
//magma_zcopy( dofs, t.dval, 1 , x->dval, 1 ); // x = t
}
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_zjacobiiter */
extern "C" magma_int_t
magma_zjacobiiter(
magma_z_sparse_matrix M, magma_z_vector c, magma_z_vector *x,
magma_z_solver_par *solver_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// local variables
magmaDoubleComplex c_zero = MAGMA_Z_ZERO, c_one = MAGMA_Z_ONE,
c_mone = MAGMA_Z_NEG_ONE;
magma_int_t dofs = M.num_rows;
magma_z_vector t, swap;
magma_z_vinit( &t, Magma_DEV, dofs, c_zero, queue );
for( magma_int_t i=0; i<solver_par->maxiter; i++ ) {
magma_z_spmv( c_mone, M, *x, c_zero, t, queue ); // t = - M * x
magma_zaxpy( dofs, c_one , c.dval, 1 , t.dval, 1 ); // t = t + c
// swap so that x again contains solution, and y is ready to be used
swap = *x;
*x = t;
t = swap;
//magma_zcopy( dofs, t.dval, 1 , x->dval, 1 ); // x = t
}
magma_z_vfree( &t, queue );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_zjacobiiter */
extern "C" magma_int_t
magma_sresidual(
magma_s_matrix A, magma_s_matrix b, magma_s_matrix x,
float *res,
magma_queue_t queue )
{
magma_int_t info = 0;
// set queue for old dense routines
magma_queue_t orig_queue=NULL;
magmablasGetKernelStream( &orig_queue );
// some useful variables
float zero = MAGMA_S_ZERO, one = MAGMA_S_ONE,
mone = MAGMA_S_NEG_ONE;
magma_int_t dofs = A.num_rows;
magma_int_t num_vecs = b.num_rows*b.num_cols/A.num_rows;
magma_s_matrix r={Magma_CSR};
if ( A.num_rows == b.num_rows ) {
CHECK( magma_svinit( &r, Magma_DEV, A.num_rows, b.num_cols, zero, queue ));
CHECK( magma_s_spmv( one, A, x, zero, r, queue )); // r = A x
magma_saxpy(dofs, mone, b.dval, 1, r.dval, 1); // r = r - b
*res = magma_snrm2(dofs, r.dval, 1); // res = ||r||
// /magma_snrm2(dofs, b.dval, 1); /||b||
//printf( "relative residual: %e\n", *res );
} else if ((b.num_rows*b.num_cols)%A.num_rows== 0 ) {
CHECK( magma_svinit( &r, Magma_DEV, b.num_rows,b.num_cols, zero, queue ));
CHECK( magma_s_spmv( one, A, x, zero, r, queue )); // r = A x
for( magma_int_t i=0; i<num_vecs; i++) {
magma_saxpy(dofs, mone, b(i), 1, r(i), 1); // r = r - b
res[i] = magma_snrm2(dofs, r(i), 1); // res = ||r||
}
// /magma_snrm2(dofs, b.dval, 1); /||b||
//printf( "relative residual: %e\n", *res );
} else {
printf("error: dimensions do not match.\n");
info = MAGMA_ERR_NOT_SUPPORTED;
}
cleanup:
magma_smfree(&r, queue );
magmablasSetKernelStream( orig_queue );
return info;
}
extern "C" magma_int_t
magma_z_applyprecond(
magma_z_sparse_matrix A,
magma_z_vector b,
magma_z_vector *x,
magma_z_preconditioner *precond,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
if ( precond->solver == Magma_JACOBI ) {
magma_zjacobi_diagscal( A.num_rows, precond->d, b, x, queue );
}
else if ( precond->solver == Magma_PASTIX ) {
magma_zapplypastix( b, x, precond, queue );
}
else if ( precond->solver == Magma_ILU ) {
magma_z_vector tmp;
magma_z_vinit( &tmp, Magma_DEV, A.num_rows, MAGMA_Z_ZERO, queue );
magma_z_vfree( &tmp, queue );
}
else if ( precond->solver == Magma_ICC ) {
magma_z_vector tmp;
magma_z_vinit( &tmp, Magma_DEV, A.num_rows, MAGMA_Z_ZERO, queue );
magma_z_vfree( &tmp, queue );
}
else if ( precond->solver == Magma_NONE ) {
magma_zcopy( b.num_rows, b.dval, 1, x->dval, 1 ); // x = b
}
else {
printf( "error: preconditioner type not yet supported.\n" );
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_svtranspose(
magma_s_matrix x,
magma_s_matrix *y,
magma_queue_t queue )
{
magma_int_t info = 0;
magma_int_t m = x.num_rows;
magma_int_t n = x.num_cols;
// set queue for old dense routines
magma_queue_t orig_queue=NULL;
magmablasGetKernelStream( &orig_queue );
magma_s_matrix x_d={Magma_CSR}, y_d={Magma_CSR};
if ( x.memory_location == Magma_DEV ) {
CHECK( magma_svinit( y, Magma_DEV, x.num_rows,x.num_cols, MAGMA_S_ZERO, queue ));
y->num_rows = x.num_rows;
y->num_cols = x.num_cols;
y->storage_type = x.storage_type;
if ( x.major == MagmaColMajor) {
y->major = MagmaRowMajor;
magmablas_stranspose( m, n, x.val, m, y->val, n );
}
else {
y->major = MagmaColMajor;
magmablas_stranspose( n, m, x.val, n, y->val, m );
}
} else {
CHECK( magma_smtransfer( x, &x_d, Magma_CPU, Magma_DEV, queue ));
CHECK( magma_svtranspose( x_d, &y_d, queue ));
CHECK( magma_smtransfer( y_d, y, Magma_DEV, Magma_CPU, queue ));
}
cleanup:
if( info != 0 ){
magma_smfree( y, queue );
}
magma_smfree( &x_d, queue );
magma_smfree( &y_d, queue );
magmablasSetKernelStream( orig_queue );
return info;
}
void magma_task_dev_dgemm(Schedule* sched_obj )
{
magma_int_t deviceID;
magma_trans_t transA;
magma_trans_t transB;
magma_int_t m;
magma_int_t n;
magma_int_t k;
double alpha;
double *dA;
magma_int_t lda;
double *dB;
magma_int_t ldb;
double beta;
double *dC;
magma_int_t ldc;
#if (dbglevel >=1)
ca_trace_start();
#endif
schedule_unpack_args_14(sched_obj, deviceID, transA, transB, m, n, k, alpha, dA, lda, dB, ldb, beta, dC, ldc);
magma_setdevice(deviceID);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(m, k, dA, lda, "A before magma_dgemm");
ca_dbg_printMat_transpose_gpu(k, n, dB, ldb, "B before magma_dgemm");
ca_dbg_printMat_transpose_gpu(m, n, dC, ldc, "C before magma_dgemm");
#endif
pthread_mutex_lock(&mutex_compute_stream);
magmablasSetKernelStream(compute_stream[deviceID]);
magma_dgemm( transA, transB, m, n, k, alpha, dA, lda, dB, ldb, beta, dC, ldc );
//task_magma_dgemm(MagmaNoTrans,MagmaNoTrans, gpu_ncols, gpu_nrows, nb, c_neg_one, dAT(K,K+A_N), dAT_LD, dAT(K+1,K), dAT_LD, c_one, dAT(K+1,K+A_N), dAT_LD);
pthread_mutex_unlock(&mutex_compute_stream);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(m, n, dC, ldc, "C after magma_dgemm");
#endif
#if (dbglevel >=1)
ca_trace_end_gpu('S');
ca_trace_end_cpu('C');
#endif
}
void magma_task_dev_dtrsm(Schedule* sched_obj )
{
magma_int_t deviceID;
magma_side_t side;
magma_uplo_t uplo;
magma_trans_t trans;
magma_diag_t diag;
magma_int_t m;
magma_int_t n;
double alpha;
double *dA;
magma_int_t lda;
double *dB;
magma_int_t ldb;
#if (dbglevel >=1)
ca_trace_start();
#endif
schedule_unpack_args_12(sched_obj, deviceID, side, uplo, trans, diag, m, n, alpha, dA, lda, dB, ldb);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(m, n, dB, ldb, "A(K,K:N) before magma_dtrsm");
#endif
magma_setdevice(deviceID);
pthread_mutex_lock(&mutex_compute_stream);
magmablasSetKernelStream(compute_stream[deviceID]);
magma_dtrsm( side, uplo, trans, diag, m, n, alpha, dA, lda, dB, ldb );
//task_magma_dtrsm('R', 'U', MagmaNoTrans, 'U', gpu_ncols, nb, c_one, dAT(K,K), dAT_LD, dAT(K,K+A_N), dAT_LD);
pthread_mutex_unlock(&mutex_compute_stream);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(m, n, dB, ldb, "A(K,K:N) after magma_dtrsm");
#endif
#if (dbglevel >=1)
ca_trace_end_gpu('U');
ca_trace_end_cpu('C');
#endif
}
void magma_task_dev_dlaswp(Schedule* sched_obj )
{
magma_int_t deviceID;
magma_int_t n;
double *dA;
magma_int_t lda;
magma_int_t i1;
magma_int_t i2;
magma_int_t *ipiv;
magma_int_t inci;
#if (dbglevel >=1)
ca_trace_start();
#endif
schedule_unpack_args_8(sched_obj, deviceID, n, dA, lda, i1, i2, ipiv, inci);
magma_setdevice(deviceID);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(n, n, dA, lda, "A(n,n) before magma_dlaswp");
#endif
pthread_mutex_lock(&mutex_compute_stream);
magmablasSetKernelStream(compute_stream[deviceID]);
magmablas_dlaswp( n, dA, lda, i1, i2, ipiv, inci );
pthread_mutex_unlock(&mutex_compute_stream);
//task_magma_dlaswp(gpu_ncols, dAT(K,K+A_N), dAT_LD, c_one, nb, &ipiv[K], c_one);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(n, n, dA, lda, "A(n,n) after magma_dlaswp");
#endif
#if (dbglevel >=1)
ca_trace_end_gpu('W');
ca_trace_end_cpu('C');
#endif
}
/**
Purpose
-------
SPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**T * U, if UPLO = MagmaUpper, or
dA = L * L**T, 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**T * U or dA = L * L**T.
@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_spotrf2_mgpu(int num_gpus, 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,
float **d_lA, magma_int_t ldda,
float **d_lP, magma_int_t lddp,
float *A, magma_int_t lda, magma_int_t h,
magma_queue_t stream[][3], magma_event_t event[][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, dd, d, 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_event_t event0[MagmaMaxGPUs], // syrk
// event1[MagmaMaxGPUs], // send off-diagonal
// event2[MagmaMaxGPUs], // send diagonal
// event3[MagmaMaxGPUs]; // trsm
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
int stream0 = 0, stream1 = 1;
#ifdef STRSM_WORK
float *d_dinvA[MagmaMaxGPUs][2], *d_x[MagmaMaxGPUs][2]; /* used by strsm_work */
#endif
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (!upper && num_gpus*ldda < max(1,n)) {
*info = -4;
} else if (upper && ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
for( d=0; d < num_gpus; d++ ) {
/* local-n and local-ld */
if (upper) {
n_local[d] = ((n/nb)/num_gpus)*nb;
if (d < (n/nb)%num_gpus)
n_local[d] += nb;
else if (d == (n/nb)%num_gpus)
n_local[d] += n%nb;
} else {
n_local[d] = ((m/nb)/num_gpus)*nb;
if (d < (m/nb)%num_gpus)
n_local[d] += nb;
else if (d == (m/nb)%num_gpus)
n_local[d] += m%nb;
}
//magma_setdevice(d);
//magma_event_create( &event0[d] );
//magma_event_create( &event1[d] );
//magma_event_create( &event2[d] );
//magma_event_create( &event3[d] );
}
magma_setdevice(0);
/* == initialize the trace */
trace_init( 1, num_gpus, 3, (magma_queue_t*)stream );
/* Use blocked code. */
if (upper) {
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
#if defined(PRECISION_d) && defined(STRSM_WORK)
/* invert the diagonals
* Allocate device memory for the inversed diagonal blocks, size=m*NB
*/
for( d=0; d < num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j < 2; j++ ) {
magma_smalloc( &d_dinvA[d][j], nb*nb );
magma_smalloc( &d_x[d][j], n*nb );
cudaMemset(d_dinvA[d][j], 0, nb*nb*sizeof(float));
cudaMemset(d_x[d][j], 0, n*nb*sizeof(float));
}
}
magma_setdevice(0);
#endif
for (j=0; j < m; j += nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (m-j));
if ( j > 0 ) {
/* needed on pluto... */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] ); // wait for the column on CPU
/* broadcast off-diagonal column to all gpus */
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; dd++ ) {
if ( d != id ) {
magma_setdevice(d);
/* wait for it on CPU */
magma_queue_wait_event( stream[d][stream0], event[id][1] );
/* send it to GPU */
trace_gpu_start( d, stream0, "comm", "rows to GPUs" );
magma_ssetmatrix_async( j, jb,
Aup(0,j), lda,
dlP(d,jb,0,buf), lddp,
stream[d][stream0] );
trace_gpu_end( d, stream0 );
magma_event_record( event[d][1], stream[d][stream0] );
}
d = (d+1)%num_gpus;
}
}
/* Update the current diagonal block */
magma_setdevice(id);
if ( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
trace_gpu_start( id, stream1, "syrk", "syrk" );
magma_ssyrk(MagmaUpper, MagmaTrans, 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 );
magma_event_record( event[id][0], stream[id][stream1] );
}
/* send the diagonal to cpu */
magma_queue_wait_event( stream[id][stream0], event[id][0] ); // wait for syrk
trace_gpu_start( id, stream0, "comm", "D to CPU" );
magma_sgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
stream[id][stream0] );
trace_gpu_end( id, stream0 );
if ( j > 0 ) {
/* Compute the local block column of the panel. */
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; dd++ ) {
j_local2 = j_local+1;
if ( d > id ) j_local2 --;
nb0 = nb*j_local2;
if ( n_local[d] > nb0 ) {
/* wait for the off-diagonal */
if ( d != id ) {
//magma_queue_sync( stream[id][3] );
dlpanel = dlP(d, jb, 0, buf);
ldpanel = lddp;
/* wait for the offdiagonal column */
magma_queue_wait_event( stream[d][stream1], event[d][1] );
} else {
dlpanel = dlA(d, 0, nb*j_local);
ldpanel = ldda;
}
/* update the panel */
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream1]);
trace_gpu_start( d, stream1, "gemm", "gemm" );
magma_sgemm(MagmaTrans, 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, stream1 );
}
d = (d+1)%num_gpus;
}
}
/* factor the diagonal */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] ); // wait for the diagonal
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 */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; dd++ ) {
magma_setdevice(d);
if ( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d, 0, 0, buf);
ldpanel = lddp;
}
trace_gpu_start( d, stream0, "comm", "D to GPUs" );
magma_ssetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, ldpanel,
stream[d][stream0] );
trace_gpu_end( d, stream0 );
magma_event_record( event[d][2], stream[d][stream0] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
trace_gpu_start( id, stream0, "comm", "D to GPUs" );
magma_ssetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
stream[id][stream0] );
trace_gpu_end( id, stream0 );
}
/* panel-factorize the off-diagonal */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; 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]-nb*j_local2;
nb0 = min(nb, nb2 );
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream1]);
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for the diagonal
if ( j+jb < m && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead the column */
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if defined(PRECISION_d) && defined(STRSM_WORK)
magmablas_strsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d][0], d_x[d][0] );
/*nb2 = n_local[d] - j_local2*nb;
magmablas_strsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d], d_x[d] ); */
#else
/*nb2 = n_local[d] - j_local2*nb;
magma_strsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldda,
dlA(d, j, nb*j_local2), ldda);
*/
magma_strsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream1 );
magma_event_record( event[d][3], stream[d][stream1] );
/* send the column to cpu */
if ( j+jb < m ) {
trace_gpu_start( d, stream0, "comm", "rows to CPU" );
magma_queue_wait_event( stream[d][stream0], event[d][3] ); // wait for lookahead
magma_sgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
stream[d][stream0] );
trace_gpu_end( d, stream0 );
magma_event_record( event[d][1], stream[d][stream0] );
}
/* update the remaining blocks */
nb2 = nb2 - nb0;
#if defined(PRECISION_d) && defined(STRSM_WORK)
magmablas_strsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_strsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda);
#endif
} else if ( nb2 > 0 ) {
/* update the entire trailing matrix */
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if defined(PRECISION_d) && defined(STRSM_WORK)
magmablas_strsm_work( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_strsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream1 );
}
d = (d+1)%num_gpus;
}
} /* end of strsm */
} /* end of for j=1, .., n */
} else {
/* -------------------------------------------- */
/* Lower-triangular case */
/* Compute the Cholesky factorization A = L*L'. */
/* -------------------------------------------- */
#if defined(PRECISION_d) && defined(STRSM_WORK)
/*
* Allocate device memory for the inversed diagonal blocks, size=N*BLOCK_SIZE
*/
for( d=0; d < num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j < 2; j++ ) {
magma_smalloc( &d_dinvA[d][j], nb*nb );
magma_smalloc( &d_x[d][j], nb*m );
cudaMemset(d_dinvA[d][j], 0, nb*nb*sizeof(float));
cudaMemset(d_x[d][j], 0, nb* m*sizeof(float));
}
}
magma_setdevice(0);
#endif
for (j=0; j < n; j += nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (n-j));
if ( j > 0 ) {
/* needed on pluto... */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] ); // wait for the column on CPU
/* broadcast offdiagonal row to all gpus */
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; dd++ ) {
if ( d != id ) {
magma_setdevice(d);
/* wait for it on CPU */
magma_queue_wait_event( stream[d][stream0], event[id][1] );
/* send it to GPU */
magma_ssetmatrix_async( jb, j,
Alo(j,0), lda,
dlPT(d,0,jb,buf), nb,
stream[d][stream0] );
magma_event_record( event[d][1], stream[d][stream0] );
}
d = (d+1)%num_gpus;
}
}
/* Update the current diagonal block */
magma_setdevice(id);
if ( j > 0 ) {
magmablasSetKernelStream(stream[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);
magma_event_record( event[id][0], stream[id][stream1] );
}
/* send the diagonal to cpu */
magma_queue_wait_event( stream[id][stream0], event[id][0] ); // wait for syrk
magma_sgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
stream[id][stream0] );
/* update the offdiagonal blocks */
if ( j > 0 ) {
/* compute the block-rows of the panel */
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; dd++ ) {
j_local2 = j_local+1;
if ( d > id ) j_local2 --;
nb0 = nb*j_local2;
if ( nb0 < n_local[d] ) {
if ( d != id ) {
dlpanel = dlPT(d, 0, jb, buf);
ldpanel = nb;
/* wait for offdiagonal row */
magma_queue_wait_event( stream[d][stream1], event[d][1] );
} else {
dlpanel = dlA(d, nb*j_local, 0);
ldpanel = ldda;
}
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream1]);
magma_sgemm( MagmaNoTrans, MagmaTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, ldpanel,
c_one, dlA(d, nb0, j), ldda);
}
d = (d+1)%num_gpus;
}
}
/* factor the diagonal */
magma_setdevice(id);
magma_queue_sync( stream[id][stream0] );
lapackf77_spotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus */
if ( (j+jb) < m ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; dd++ ) {
magma_setdevice(d);
if ( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
magma_ssetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, ldpanel,
stream[d][stream0] );
magma_event_record( event[d][2], stream[d][stream0] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
magma_ssetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
stream[id][stream0] );
}
/* factorize off-diagonal blocks */
if ( (j+jb) < m ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd < num_gpus; 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);
magmablasSetKernelStream(stream[d][stream1]);
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for the diagonal
if ( j+jb < n && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead the column */
#if defined(PRECISION_d) && defined(STRSM_WORK)
magmablas_strsm_work( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
d_dinvA[d][0], d_x[d][0]);
#else
magma_strsm( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
magma_event_record( event[d][3], stream[d][stream1] );
/* send the column to cpu */
if ( j+jb < n ) {
magma_queue_wait_event( stream[d][stream0], event[d][3] ); // wait for lookahead
magma_sgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
stream[d][stream0] );
magma_event_record( event[d][1], stream[d][stream0] );
}
/* update the remaining blocks */
nb2 = nb2 - nb0;
#if defined(PRECISION_d) && defined(STRSM_WORK)
magmablas_strsm_work( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_strsm( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda);
#endif
} else if ( nb2 > 0 ) {
/* update the entire trailing matrix */
#if defined(PRECISION_d) && defined(STRSM_WORK)
magmablas_strsm_work( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
d_dinvA[d][1], d_x[d][1] );
#else
magma_strsm( MagmaRight, MagmaLower, MagmaTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
}
d = (d+1)%num_gpus;
}
}
}
} /* end of else not upper */
/* == finalize the trace == */
trace_finalize( "spotrf.svg", "trace.css" );
/* clean up */
for( d=0; d < num_gpus; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
//magma_event_destroy( event0[d] );
//magma_event_destroy( event1[d] );
//magma_event_destroy( event2[d] );
//magma_event_destroy( event3[d] );
}
magma_setdevice(0);
return *info;
} /* magma_spotrf_mgpu */
/**
Purpose
=======
SSYTRF_nopiv_gpu computes the LDLt factorization of a real symmetric
matrix A.
The factorization has the form
A = U^H * D * U , if UPLO = 'U', or
A = L * D * L^H, if UPLO = 'L',
where U is an upper triangular matrix, L is lower triangular, and
D is a diagonal matrix.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
UPLO CHARACTER*1
- = 'U': Upper triangle of A is stored;
- = 'L': 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 (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U^H D U or A = L D L^H.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using cudaMallocHost.
@param[in]
LDA INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
INFO INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -6, the GPU memory allocation failed
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_ssytrf_comp
******************************************************************* */
extern "C" magma_int_t
magma_ssytrf_nopiv_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaFloat_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define A(i, j) (A)
#define dA(i, j) (dA +(j)*ldda + (i))
#define dW(i, j) (dW +(j)*ldda + (i))
#define dWt(i, j) (dW +(j)*nb + (i))
/* Local variables */
float zone = MAGMA_S_ONE;
float mzone = MAGMA_S_NEG_ONE;
int upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, nb, ib, iinfo;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
/* Quick return */
if ( n == 0 )
return MAGMA_SUCCESS;
nb = magma_get_ssytrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_t stream[2];
magma_event_t event;
magma_queue_create(&stream[0]);
magma_queue_create(&stream[1]);
magma_event_create( &event );
trace_init( 1, 1, 2, stream );
// CPU workspace
float *A;
if (MAGMA_SUCCESS != magma_smalloc_pinned( &A, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// GPU workspace
magmaFloat_ptr dW;
if (MAGMA_SUCCESS != magma_smalloc( &dW, (1+nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// main loop
for (j=0; j<n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
//magma_queue_wait_event( stream[1], event );
magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), nb, stream[1]);
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(stream[1]);
trace_cpu_start( 0, "potrf", "potrf" );
ssytrf_nopiv_cpu(MagmaUpper, jb, ib, A(j, j), nb, info);
trace_cpu_end( 0 );
if (*info != 0){
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(jb, jb, A(j, j), nb, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaUnit,
jb, (n-j-jb),
zone, dA(j, j), ldda,
dA(j, j+jb), ldda);
magma_scopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb );
// update the trailing submatrix with D
magmablas_slascl_diag(MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
&iinfo);
trace_gpu_end( 0, 0 );
// update the trailing submatrix with U and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k<n; k+=nb) {
magma_int_t kb = min(nb,n-k);
magma_sgemm(MagmaConjTrans, MagmaNoTrans, kb, n-k, jb,
mzone, dWt(0, k), nb,
dA(j, k), ldda,
zone, dA(k, k), ldda);
if (k==j+jb)
magma_event_record( event, stream[0] );
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// main loop
for (j=0; j<n; j+=nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
//magma_queue_wait_event( stream[0], event );
magma_event_sync(event);
magma_sgetmatrix_async(jb, jb, dA(j, j), ldda, A(j,j), nb, stream[1]);
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync(stream[1]);
trace_cpu_start( 0, "potrf", "potrf" );
ssytrf_nopiv_cpu(MagmaLower, jb, ib, A(j, j), nb, info);
trace_cpu_end( 0 );
if (*info != 0){
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_ssetmatrix_async(jb, jb, A(j, j), nb, dA(j, j), ldda, stream[0]);
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
magmablasSetKernelStream( stream[0] );
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_strsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaUnit,
(n-j-jb), jb,
zone, dA(j, j), ldda,
dA(j+jb, j), ldda);
magma_scopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda );
// update the trailing submatrix with D
magmablas_slascl_diag(MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
&iinfo);
trace_gpu_end( 0, 0 );
// update the trailing submatrix with L and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k<n; k+=nb) {
magma_int_t kb = min(nb,n-k);
magma_sgemm(MagmaNoTrans, MagmaConjTrans, n-k, kb, jb,
mzone, dA(k, j), ldda,
dW(k, 0), ldda,
zone, dA(k, k), ldda);
if (k==j+jb)
magma_event_record( event, stream[0] );
}
trace_gpu_end( 0, 0 );
}
}
}
trace_finalize( "ssytrf.svg","trace.css" );
magma_queue_destroy(stream[0]);
magma_queue_destroy(stream[1]);
magma_event_destroy( event );
magma_free( dW );
magma_free_pinned( A );
magmablasSetKernelStream( orig_stream );
return MAGMA_SUCCESS;
} /* magma_ssytrf_nopiv */
/**
Purpose
-------
SLAEX3 finds the roots of the secular equation, as defined by the
values in D, W, and RHO, between 1 and K. It makes the
appropriate calls to SLAED4 and then updates the eigenvectors by
multiplying the matrix of eigenvectors of the pair of eigensystems
being combined by the matrix of eigenvectors of the K-by-K system
which is solved here.
It is used in the last step when only a part of the eigenvectors
is required.
It compute only the required part of the eigenvectors and the rest
is not used.
This code makes very mild assumptions about floating point
arithmetic. It will work on machines with a guard digit in
add/subtract, or on those binary machines without guard digits
which subtract like the Cray X-MP, Cray Y-MP, Cray C-90, or Cray-2.
It could conceivably fail on hexadecimal or decimal machines
without guard digits, but we know of none.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
k INTEGER
The number of terms in the rational function to be solved by
SLAED4. K >= 0.
@param[in]
n INTEGER
The number of rows and columns in the Q matrix.
N >= K (deflation may result in N > K).
@param[in]
n1 INTEGER
The location of the last eigenvalue in the leading submatrix.
min(1,N) <= N1 <= N/2.
@param[out]
d REAL array, dimension (N)
D(I) contains the updated eigenvalues for
1 <= I <= K.
@param[out]
Q REAL array, dimension (LDQ,N)
Initially the first K columns are used as workspace.
On output the columns ??? to ??? contain
the updated eigenvectors.
@param[in]
ldq INTEGER
The leading dimension of the array Q. LDQ >= max(1,N).
@param[in]
rho REAL
The value of the parameter in the rank one update equation.
RHO >= 0 required.
@param[in,out]
dlamda REAL array, dimension (K)
The first K elements of this array contain the old roots
of the deflated updating problem. These are the poles
of the secular equation. May be changed on output by
having lowest order bit set to zero on Cray X-MP, Cray Y-MP,
Cray-2, or Cray C-90, as described above.
@param[in]
Q2 REAL array, dimension (LDQ2, N)
The first K columns of this matrix contain the non-deflated
eigenvectors for the split problem.
@param[in]
indx INTEGER array, dimension (N)
The permutation used to arrange the columns of the deflated
Q matrix into three groups (see SLAED2).
The rows of the eigenvectors found by SLAED4 must be likewise
permuted before the matrix multiply can take place.
@param[in]
ctot INTEGER array, dimension (4)
A count of the total number of the various types of columns
in Q, as described in INDX. The fourth column type is any
column which has been deflated.
@param[in,out]
w REAL array, dimension (K)
The first K elements of this array contain the components
of the deflation-adjusted updating vector. Destroyed on
output.
@param
s (workspace) REAL array, dimension (N1 + 1)*K
Will contain the eigenvectors of the repaired matrix which
will be multiplied by the previously accumulated eigenvectors
to update the system.
@param[out]
indxq INTEGER array, dimension (N)
On exit, the permutation which will reintegrate the
subproblems back into sorted order,
i.e. D( INDXQ( I = 1, N ) ) will be in ascending order.
@param
dwork (devices workspaces) REAL array of arrays,
dimension NRGPU.
if NRGPU = 1 the dimension of the first workspace
should be (3*N*N/2+3*N)
otherwise the NRGPU workspaces should have the size
ceil((N-N1) * (N-N1) / floor(ngpu/2)) +
NB * ((N-N1) + (N-N1) / floor(ngpu/2))
@param
queues (device queues) magma_queue_t array,
dimension (MagmaMaxGPUs,2)
@param[in]
range magma_range_t
- = MagmaRangeAll: all eigenvalues will be found.
- = MagmaRangeV: all eigenvalues in the half-open interval (VL,VU]
will be found.
- = MagmaRangeI: the IL-th through IU-th eigenvalues will be found.
TODO verify range, vl, vu, il, iu -- copied from slaex1.
@param[in]
vl REAL
@param[in]
vu REAL
if RANGE=MagmaRangeV, the lower and upper bounds of the interval to
be searched for eigenvalues. VL < VU.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeI.
@param[in]
il INTEGER
@param[in]
iu INTEGER
if RANGE=MagmaRangeI, the indices (in ascending order) of the
smallest and largest eigenvalues to be returned.
1 <= IL <= IU <= N, if N > 0; IL = 1 and IU = 0 if N = 0.
Not referenced if RANGE = MagmaRangeAll or MagmaRangeV.
@param[out]
info INTEGER
- = 0: successful exit.
- < 0: if INFO = -i, the i-th argument had an illegal value.
- > 0: if INFO = 1, an eigenvalue did not converge
Further Details
---------------
Based on contributions by
Jeff Rutter, Computer Science Division, University of California
at Berkeley, USA
Modified by Francoise Tisseur, University of Tennessee.
@ingroup magma_ssyev_aux
********************************************************************/
extern "C" magma_int_t
magma_slaex3_m(
magma_int_t ngpu,
magma_int_t k, magma_int_t n, magma_int_t n1, float *d,
float *Q, magma_int_t ldq, float rho,
float *dlamda, float *Q2, magma_int_t *indx,
magma_int_t *ctot, float *w, float *s, magma_int_t *indxq,
magmaFloat_ptr dwork[],
magma_queue_t queues[MagmaMaxGPUs][2],
magma_range_t range, float vl, float vu, magma_int_t il, magma_int_t iu,
magma_int_t *info )
{
#define Q(i_,j_) (Q + (i_) + (j_)*ldq)
#define dQ2(id) (dwork[id])
#define dS(id, ii) (dwork[id] + n2*n2_loc + (ii)*(n2*nb))
#define dQ(id, ii) (dwork[id] + n2*n2_loc + 2*(n2*nb) + (ii)*(n2_loc*nb))
if (ngpu == 1) {
magma_setdevice(0);
magma_slaex3(k, n, n1, d, Q, ldq, rho,
dlamda, Q2, indx, ctot, w, s, indxq,
*dwork, range, vl, vu, il, iu, info );
return *info;
}
float d_one = 1.;
float d_zero = 0.;
magma_int_t ione = 1;
magma_int_t ineg_one = -1;
magma_int_t iil, iiu, rk;
magma_int_t n1_loc, n2_loc, ib, nb, ib2, igpu;
magma_int_t ni_loc[MagmaMaxGPUs];
magma_int_t i, ind, iq2, j, n12, n2, n23, tmp;
float temp;
magma_int_t alleig, valeig, indeig;
alleig = (range == MagmaRangeAll);
valeig = (range == MagmaRangeV);
indeig = (range == MagmaRangeI);
*info = 0;
if (k < 0)
*info=-1;
else if (n < k)
*info=-2;
else if (ldq < max(1,n))
*info=-6;
else if (! (alleig || valeig || indeig))
*info = -15;
else {
if (valeig) {
if (n > 0 && vu <= vl)
*info = -17;
}
else if (indeig) {
if (il < 1 || il > max(1,n))
*info = -18;
else if (iu < min(n,il) || iu > n)
*info = -19;
}
}
if (*info != 0) {
magma_xerbla(__func__, -(*info));
return *info;
}
// Quick return if possible
if (k == 0)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/*
Modify values DLAMDA(i) to make sure all DLAMDA(i)-DLAMDA(j) can
be computed with high relative accuracy (barring over/underflow).
This is a problem on machines without a guard digit in
add/subtract (Cray XMP, Cray YMP, Cray C 90 and Cray 2).
The following code replaces DLAMDA(I) by 2*DLAMDA(I)-DLAMDA(I),
which on any of these machines zeros out the bottommost
bit of DLAMDA(I) if it is 1; this makes the subsequent
subtractions DLAMDA(I)-DLAMDA(J) unproblematic when cancellation
occurs. On binary machines with a guard digit (almost all
machines) it does not change DLAMDA(I) at all. On hexadecimal
and decimal machines with a guard digit, it slightly
changes the bottommost bits of DLAMDA(I). It does not account
for hexadecimal or decimal machines without guard digits
(we know of none). We use a subroutine call to compute
2*DLAMBDA(I) to prevent optimizing compilers from eliminating
this code.*/
//#define CHECK_CPU
#ifdef CHECK_CPU
float *hwS[2][MagmaMaxGPUs], *hwQ[2][MagmaMaxGPUs], *hwQ2[MagmaMaxGPUs];
#define hQ2(id) (hwQ2[id])
#define hS(id, ii) (hwS[ii][id])
#define hQ(id, ii) (hwQ[ii][id])
#endif
n2 = n - n1;
n12 = ctot[0] + ctot[1];
n23 = ctot[1] + ctot[2];
iq2 = n1 * n12;
//lq2 = iq2 + n2 * n23;
n1_loc = (n1-1) / (ngpu/2) + 1;
n2_loc = (n2-1) / (ngpu/2) + 1;
nb = magma_get_slaex3_m_nb();
if (n1 >= magma_get_slaex3_m_k()) {
#ifdef CHECK_CPU
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_smalloc_pinned( &(hwS[0][igpu]), n2*nb );
magma_smalloc_pinned( &(hwS[1][igpu]), n2*nb );
magma_smalloc_pinned( &(hwQ2[igpu]), n2*n2_loc );
magma_smalloc_pinned( &(hwQ[0][igpu]), n2_loc*nb );
magma_smalloc_pinned( &(hwQ[1][igpu]), n2_loc*nb );
}
#endif
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
ni_loc[igpu] = min(n1_loc, n1 - igpu/2 * n1_loc);
#ifdef CHECK_CPU
lapackf77_slacpy("A", &ni_loc[igpu], &n12, Q2+n1_loc*(igpu/2), &n1, hQ2(igpu), &n1_loc);
#endif
magma_setdevice(igpu);
magma_ssetmatrix_async( ni_loc[igpu], n12,
Q2+n1_loc*(igpu/2), n1,
dQ2(igpu), n1_loc, queues[igpu][0] );
ni_loc[igpu+1] = min(n2_loc, n2 - igpu/2 * n2_loc);
#ifdef CHECK_CPU
lapackf77_slacpy("A", &ni_loc[igpu+1], &n23, Q2+iq2+n2_loc*(igpu/2), &n2, hQ2(igpu+1), &n2_loc);
#endif
magma_setdevice(igpu+1);
magma_ssetmatrix_async( ni_loc[igpu+1], n23,
Q2+iq2+n2_loc*(igpu/2), n2,
dQ2(igpu+1), n2_loc, queues[igpu+1][0] );
}
}
//
#ifdef _OPENMP
/////////////////////////////////////////////////////////////////////////////////
//openmp implementation
/////////////////////////////////////////////////////////////////////////////////
magma_timer_t time=0;
timer_start( time );
#pragma omp parallel private(i, j, tmp, temp)
{
magma_int_t id = omp_get_thread_num();
magma_int_t tot = omp_get_num_threads();
magma_int_t ib = ( id * k) / tot; //start index of local loop
magma_int_t ie = ((id+1) * k) / tot; //end index of local loop
magma_int_t ik = ie - ib; //number of local indices
for (i = ib; i < ie; ++i)
dlamda[i]=lapackf77_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = ib; j < ie; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&k, &tmpp, dlamda, w, Q(0,j), &rho, &d[j], &iinfo);
// If the zero finder fails, the computation is terminated.
if (iinfo != 0) {
#pragma omp critical (info)
*info = iinfo;
break;
}
}
#pragma omp barrier
if (*info == 0) {
#pragma omp single
{
//Prepare the INDXQ sorting permutation.
magma_int_t nk = n - k;
lapackf77_slamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(k, indxq, &iil, &iiu, il, iu);
else {
iil = 1;
iiu = k;
}
rk = iiu - iil + 1;
}
if (k == 2) {
#pragma omp single
{
for (j = 0; j < k; ++j) {
w[0] = *Q(0,j);
w[1] = *Q(1,j);
i = indx[0] - 1;
*Q(0,j) = w[i];
i = indx[1] - 1;
*Q(1,j) = w[i];
}
}
}
else if (k != 1) {
// Compute updated W.
blasf77_scopy( &ik, &w[ib], &ione, &s[ib], &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &ik, Q(ib,ib), &tmp, &w[ib], &ione);
for (j = 0; j < k; ++j) {
magma_int_t i_tmp = min(j, ie);
for (i = ib; i < i_tmp; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
i_tmp = max(j+1, ib);
for (i = i_tmp; i < ie; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
}
for (i = ib; i < ie; ++i)
w[i] = copysign( sqrt( -w[i] ), s[i]);
#pragma omp barrier
//reduce the number of used threads to have enough S workspace
tot = min(n1, omp_get_num_threads());
if (id < tot) {
ib = ( id * rk) / tot + iil - 1;
ie = ((id+1) * rk) / tot + iil - 1;
ik = ie - ib;
}
else {
ib = -1;
ie = -1;
ik = -1;
}
// Compute eigenvectors of the modified rank-1 modification.
for (j = ib; j < ie; ++j) {
for (i = 0; i < k; ++i)
s[id*k + i] = w[i] / *Q(i,j);
temp = magma_cblas_snrm2( k, s+id*k, 1 );
for (i = 0; i < k; ++i) {
magma_int_t iii = indx[i] - 1;
*Q(i,j) = s[id*k + iii] / temp;
}
}
}
}
}
if (*info != 0)
return *info;
timer_stop( time );
timer_printf( "eigenvalues/vector D+zzT = %6.2f\n", time );
#else
/////////////////////////////////////////////////////////////////////////////////
// Non openmp implementation
/////////////////////////////////////////////////////////////////////////////////
magma_timer_t time=0;
timer_start( time );
for (i = 0; i < k; ++i)
dlamda[i]=lapackf77_slamc3(&dlamda[i], &dlamda[i]) - dlamda[i];
for (j = 0; j < k; ++j) {
magma_int_t tmpp=j+1;
magma_int_t iinfo = 0;
lapackf77_slaed4(&k, &tmpp, dlamda, w, Q(0,j), &rho, &d[j], &iinfo);
// If the zero finder fails, the computation is terminated.
if (iinfo != 0)
*info=iinfo;
}
if (*info != 0)
return *info;
//Prepare the INDXQ sorting permutation.
magma_int_t nk = n - k;
lapackf77_slamrg( &k, &nk, d, &ione, &ineg_one, indxq);
//compute the lower and upper bound of the non-deflated eigenvectors
if (valeig)
magma_svrange(k, d, &iil, &iiu, vl, vu);
else if (indeig)
magma_sirange(k, indxq, &iil, &iiu, il, iu);
else {
iil = 1;
iiu = k;
}
rk = iiu - iil + 1;
if (k == 2) {
for (j = 0; j < k; ++j) {
w[0] = *Q(0,j);
w[1] = *Q(1,j);
i = indx[0] - 1;
*Q(0,j) = w[i];
i = indx[1] - 1;
*Q(1,j) = w[i];
}
}
else if (k != 1) {
// Compute updated W.
blasf77_scopy( &k, w, &ione, s, &ione);
// Initialize W(I) = Q(I,I)
tmp = ldq + 1;
blasf77_scopy( &k, Q, &tmp, w, &ione);
for (j = 0; j < k; ++j) {
for (i = 0; i < j; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
for (i = j+1; i < k; ++i)
w[i] = w[i] * ( *Q(i, j) / ( dlamda[i] - dlamda[j] ) );
}
for (i = 0; i < k; ++i)
w[i] = copysign( sqrt( -w[i] ), s[i]);
// Compute eigenvectors of the modified rank-1 modification.
for (j = iil-1; j < iiu; ++j) {
for (i = 0; i < k; ++i)
s[i] = w[i] / *Q(i,j);
temp = magma_cblas_snrm2( k, s, 1 );
for (i = 0; i < k; ++i) {
magma_int_t iii = indx[i] - 1;
*Q(i,j) = s[iii] / temp;
}
}
}
timer_stop( time );
timer_printf( "eigenvalues/vector D+zzT = %6.2f\n", time );
#endif //_OPENMP
// Compute the updated eigenvectors.
timer_start( time );
if (rk > 0) {
if (n1 < magma_get_slaex3_m_k()) {
// stay on the CPU
if ( n23 != 0 ) {
lapackf77_slacpy("A", &n23, &rk, Q(ctot[0],iil-1), &ldq, s, &n23);
blasf77_sgemm("N", "N", &n2, &rk, &n23, &d_one, &Q2[iq2], &n2,
s, &n23, &d_zero, Q(n1,iil-1), &ldq );
}
else
lapackf77_slaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if ( n12 != 0 ) {
lapackf77_slacpy("A", &n12, &rk, Q(0,iil-1), &ldq, s, &n12);
blasf77_sgemm("N", "N", &n1, &rk, &n12, &d_one, Q2, &n1,
s, &n12, &d_zero, Q(0,iil-1), &ldq);
}
else
lapackf77_slaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
else {
//use the gpus
ib = min(nb, rk);
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_ssetmatrix_async( n23, ib,
Q(ctot[0],iil-1), ldq,
dS(igpu+1,0), n23, queues[igpu+1][0] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_ssetmatrix_async( n12, ib,
Q(0,iil-1), ldq,
dS(igpu,0), n12, queues[igpu][0] );
}
}
for (i = 0; i < rk; i += nb) {
ib = min(nb, rk - i);
ind = (i/nb)%2;
if (i+nb < rk) {
ib2 = min(nb, rk - i - nb);
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_ssetmatrix_async( n23, ib2,
Q(ctot[0],iil-1+i+nb), ldq,
dS(igpu+1,(ind+1)%2), n23, queues[igpu+1][(ind+1)%2] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_ssetmatrix_async( n12, ib2,
Q(0,iil-1+i+nb), ldq,
dS(igpu,(ind+1)%2), n12, queues[igpu][(ind+1)%2] );
}
}
}
// Ensure that the data is copied on gpu since we will overwrite it.
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
#ifdef CHECK_CPU
lapackf77_slacpy("A", &n23, &ib, Q(ctot[0],iil-1+i), &ldq, hS(igpu+1,ind), &n23);
#endif
magma_setdevice(igpu+1);
magma_queue_sync( queues[igpu+1][ind] );
}
if (n12 != 0) {
#ifdef CHECK_CPU
lapackf77_slacpy("A", &n12, &ib, Q(0,iil-1+i), &ldq, hS(igpu,ind), &n12);
#endif
magma_setdevice(igpu);
magma_queue_sync( queues[igpu][ind] );
}
}
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
#ifdef CHECK_CPU
blasf77_sgemm("N", "N", &ni_loc[igpu+1], &ib, &n23, &d_one, hQ2(igpu+1), &n2_loc,
hS(igpu+1,ind), &n23, &d_zero, hQ(igpu+1, ind), &n2_loc);
#endif
magma_setdevice(igpu+1);
magmablasSetKernelStream(queues[igpu+1][ind]);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, ni_loc[igpu+1], ib, n23, d_one, dQ2(igpu+1), n2_loc,
dS(igpu+1, ind), n23, d_zero, dQ(igpu+1, ind), n2_loc);
#ifdef CHECK_CPU
printf("norm Q %d: %f\n", igpu+1, cpu_gpu_sdiff(ni_loc[igpu+1], ib, hQ(igpu+1, ind), n2_loc, dQ(igpu+1, ind), n2_loc));
#endif
}
if (n12 != 0) {
#ifdef CHECK_CPU
blasf77_sgemm("N", "N", &ni_loc[igpu], &ib, &n12, &d_one, hQ2(igpu), &n1_loc,
hS(igpu,ind%2), &n12, &d_zero, hQ(igpu, ind%2), &n1_loc);
#endif
magma_setdevice(igpu);
magmablasSetKernelStream(queues[igpu][ind]);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, ni_loc[igpu], ib, n12, d_one, dQ2(igpu), n1_loc,
dS(igpu, ind), n12, d_zero, dQ(igpu, ind), n1_loc);
#ifdef CHECK_CPU
printf("norm Q %d: %f\n", igpu, cpu_gpu_sdiff(ni_loc[igpu], ib, hQ(igpu, ind), n1_loc, dQ(igpu, ind), n1_loc));
#endif
}
}
for (igpu = 0; igpu < ngpu-1; igpu += 2) {
if (n23 != 0) {
magma_setdevice(igpu+1);
magma_sgetmatrix( ni_loc[igpu+1], ib, dQ(igpu+1, ind), n2_loc,
Q(n1+n2_loc*(igpu/2),iil-1+i), ldq );
// magma_sgetmatrix_async( ni_loc[igpu+1], ib, dQ(igpu+1, ind), n2_loc,
// Q(n1+n2_loc*(igpu/2),iil-1+i), ldq, queues[igpu+1][ind] );
}
if (n12 != 0) {
magma_setdevice(igpu);
magma_sgetmatrix( ni_loc[igpu], ib, dQ(igpu, ind), n1_loc,
Q(n1_loc*(igpu/2),iil-1+i), ldq );
// magma_sgetmatrix_async( ni_loc[igpu], ib, dQ(igpu, ind), n1_loc,
// Q(n1_loc*(igpu/2),iil-1+i), ldq, queues[igpu][ind] );
}
}
}
for (igpu = 0; igpu < ngpu; ++igpu) {
#ifdef CHECK_CPU
magma_free_pinned( hwS[1][igpu] );
magma_free_pinned( hwS[0][igpu] );
magma_free_pinned( hwQ2[igpu] );
magma_free_pinned( hwQ[1][igpu] );
magma_free_pinned( hwQ[0][igpu] );
#endif
magma_setdevice(igpu);
magma_queue_sync( queues[igpu][0] );
magma_queue_sync( queues[igpu][1] );
}
if ( n23 == 0 )
lapackf77_slaset("A", &n2, &rk, &d_zero, &d_zero, Q(n1,iil-1), &ldq);
if ( n12 == 0 )
lapackf77_slaset("A", &n1, &rk, &d_zero, &d_zero, Q(0,iil-1), &ldq);
}
}
timer_stop( time );
timer_printf( "gemms = %6.2f\n", time );
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_slaed3_m */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ssymmetrize
Code is very similar to testing_stranspose.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;
magma_int_t N, size, lda, ldda;
magma_int_t ione = 1;
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
printf("%% uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf("%% N CPU GByte/s (ms) GPU GByte/s (ms) check\n");
printf("%%====================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
size = lda*N;
// load strictly lower triangle, save strictly upper triangle
gbytes = sizeof(float) * 1.*N*(N-1) / 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 < N; ++i ) {
h_A[i + j*lda] = MAGMA_S_MAKE( i + j/10000., j );
}
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_ssetmatrix( N, N, h_A, lda, d_A, ldda );
magmablasSetKernelStream( opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
//magmablas_ssymmetrize( opts.uplo, N-2, d_A+1+ldda, ldda ); // inset by 1 row & col
magmablas_ssymmetrize( opts.uplo, N, d_A, ldda );
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gbytes / gpu_time;
/* =====================================================================
Performs operation using naive in-place algorithm
(LAPACK doesn't implement symmetrize)
=================================================================== */
cpu_time = magma_wtime();
//for( int j = 1; j < N-1; ++j ) { // inset by 1 row & col
// for( int i = 1; i < j; ++i ) {
for( int j = 0; j < N; ++j ) {
for( int i = 0; i < j; ++i ) {
if ( opts.uplo == MagmaLower ) {
h_A[i + j*lda] = MAGMA_S_CONJ( h_A[j + i*lda] );
}
else {
h_A[j + i*lda] = MAGMA_S_CONJ( h_A[i + j*lda] );
}
}
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gbytes / cpu_time;
/* =====================================================================
Check the result
=================================================================== */
magma_sgetmatrix( N, N, d_A, ldda, h_R, lda );
blasf77_saxpy(&size, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_slange("f", &N, &N, h_R, &lda, work);
printf("%5d %7.2f (%7.2f) %7.2f (%7.2f) %s\n",
(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" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CGETRF_NOPIV_GPU computes an LU factorization of a general M-by-N
matrix A without any pivoting.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_nopiv_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define dA(i,j) (dA + (i)*nb + (j)*nb*ldda)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, mindim;
magma_int_t i, rows, s, lddwork;
magmaFloatComplex *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
magma_cgetrf_nopiv( m, n, work, m, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( i=0; i < s; i++ ) {
// download i-th panel
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dA(i,i), ldda, work, lddwork, stream[0] );
if ( i > 0 ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n - (i+1)*nb,
c_one, dA(i-1,i-1), ldda,
dA(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-i*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i, i-1), ldda, dA(i-1,i+1), ldda,
c_one, dA(i, i+1), ldda );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
magma_cgetrf_nopiv( rows, nb, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dA(i, i), ldda, stream[0] );
magma_queue_sync( stream[0] );
// do the small non-parallel computations
if ( s > (i+1) ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n-s*nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
magma_cgetmatrix( rows, nb0, dA(s,s), ldda, work, lddwork );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
magma_cgetrf_nopiv( rows, nb0, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dA(s,s), ldda );
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb0, n-s*nb-nb0,
c_one, dA(s,s), ldda,
dA(s,s)+nb0, ldda);
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_cgetrf_nopiv_gpu */
/**
Purpose
-------
ZGETRF_m computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine. The matrix may exceed the GPU memory.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Note: The factorization of big panel is done calling multiple-gpu-interface.
Pivots are applied on GPU within the big panel.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX_16 array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrf_m(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaDoubleComplex *A, magma_int_t lda, magma_int_t *ipiv,
magma_int_t *info)
{
#define A(i,j) (A + (j)*lda + (i))
#define dAT(d,i,j) (dAT[d] + (i)*nb*ldn_local + (j)*nb)
#define dPT(d,i,j) (dPT[d] + (i)*nb*nb + (j)*nb*maxm)
magma_timer_t time=0, time_total=0, time_alloc=0, time_set=0, time_get=0, time_comp=0;
timer_start( time_total );
real_Double_t flops;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex *dAT[MagmaMaxGPUs], *dA[MagmaMaxGPUs], *dPT[MagmaMaxGPUs];
magma_int_t iinfo = 0, nb, nbi, maxm, n_local[MagmaMaxGPUs], ldn_local;
magma_int_t N, M, NB, NBk, I, d, ngpu0 = ngpu;
magma_int_t ii, jj, h, offset, ib, rows;
magma_queue_t stream[MagmaMaxGPUs][2];
magma_event_t event[MagmaMaxGPUs][2];
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* initialize nb */
nb = magma_get_zgetrf_nb(m);
maxm = ((m + 31)/32)*32;
/* figure out NB */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaDoubleComplex);
/* number of columns in the big panel */
h = 1+(2+ngpu0);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
//NB = 5*max(nb,32);
if ( ngpu0 > ceil((double)NB/nb) ) {
ngpu = (int)ceil((double)NB/nb);
h = 1+(2+ngpu);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
} else {
ngpu = ngpu0;
}
if ( ngpu*NB >= n ) {
#ifdef CHECK_ZGETRF_OOC
printf( " * still fit in GPU memory.\n" );
#endif
NB = n;
} else {
#ifdef CHECK_ZGETRF_OOC
printf( " * don't fit in GPU memory.\n" );
#endif
NB = ngpu*NB;
NB = max( nb, (NB / nb) * nb); /* making sure it's devisable by nb (x64) */
}
#ifdef CHECK_ZGETRF_OOC
if ( NB != n ) printf( " * running in out-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n", n, NB, nb, (double)freeMem );
else printf( " * running in in-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n", n, NB, nb, (double)freeMem );
#endif
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code for scalar of one tile. */
lapackf77_zgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
/* allocate memory on GPU to store the big panel */
timer_start( time_alloc );
n_local[0] = (NB/nb)/ngpu;
if ( NB%(nb*ngpu) != 0 )
n_local[0]++;
n_local[0] *= nb;
ldn_local = ((n_local[0]+31)/32)*32;
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_zmalloc( &dA[d], (ldn_local+h*nb)*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dPT[d] = dA[d] + nb*maxm; /* for storing the previous panel from CPU */
dAT[d] = dA[d] + h*nb*maxm; /* for storing the big panel */
magma_queue_create( &stream[d][0] );
magma_queue_create( &stream[d][1] );
magma_event_create( &event[d][0] );
magma_event_create( &event[d][1] );
}
//magma_setdevice(0);
timer_stop( time_alloc );
for( I=0; I < n; I += NB ) {
M = m;
N = min( NB, n-I ); /* number of columns in this big panel */
//s = min( max(m-I,0), N )/nb; /* number of small block-columns in this big panel */
maxm = ((M + 31)/32)*32;
if ( ngpu0 > ceil((double)N/nb) ) {
ngpu = (int)ceil((double)N/nb);
} else {
ngpu = ngpu0;
}
for( d=0; d < ngpu; d++ ) {
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;
}
ldn_local = ((n_local[0]+31)/32)*32;
/* upload the next big panel into GPU, transpose (A->A'), and pivot it */
timer_start( time );
magmablas_zsetmatrix_transpose_mgpu(ngpu, stream, A(0,I), lda,
dAT, ldn_local, dA, maxm, M, N, nb);
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
time_set += timer_stop( time );
timer_start( time );
/* == --------------------------------------------------------------- == */
/* == loop around the previous big-panels to update the new big-panel == */
for( offset = 0; offset < min(m,I); offset += NB ) {
NBk = min( m-offset, NB );
/* start sending the first tile from the previous big-panels to gpus */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
nbi = min( nb, NBk );
magma_zsetmatrix_async( (M-offset), nbi,
A(offset,offset), lda,
dA[d], (maxm-offset), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][0] );
/* transpose */
magmablas_ztranspose( M-offset, nbi, dA[d], maxm-offset, dPT(d,0,0), nb );
}
/* applying the pivot from the previous big-panel */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magmablas_zlaswp_q( ldn_local, dAT(d,0,0), ldn_local, offset+1, offset+NBk, ipiv, 1, stream[d][1] );
}
/* == going through each block-column of previous big-panels == */
for( jj=0, ib=offset/nb; jj < NBk; jj += nb, ib++ ) {
ii = offset+jj;
rows = maxm - ii;
nbi = min( nb, NBk-jj );
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
/* wait for a block-column on GPU */
magma_queue_sync( stream[d][0] );
/* start sending next column */
if ( jj+nb < NBk ) {
magma_zsetmatrix_async( (M-ii-nb), min(nb,NBk-jj-nb),
A(ii+nb,ii+nb), lda,
dA[d], (rows-nb), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][(1+jj/nb)%2] );
/* transpose next column */
magmablas_ztranspose( M-ii-nb, nb, dA[d], rows-nb, dPT(d,0,(1+jj/nb)%2), nb );
}
/* update with the block column */
magmablasSetKernelStream(stream[d][1]);
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n_local[d], nbi, c_one, dPT(d,0,(jj/nb)%2), nb, dAT(d,ib,0), ldn_local );
if ( M > ii+nb ) {
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n_local[d], M-(ii+nb), nbi, c_neg_one, dAT(d,ib,0), ldn_local,
dPT(d,1,(jj/nb)%2), nb, c_one, dAT(d,ib+1,0), ldn_local );
}
magma_event_record( event[d][(jj/nb)%2], stream[d][1] );
} /* end of for each block-columns in a big-panel */
}
} /* end of for each previous big-panels */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
/* calling magma-gpu interface to panel-factorize the big panel */
if ( M > I ) {
magma_zgetrf2_mgpu(ngpu, M-I, N, nb, I, dAT, ldn_local, ipiv+I, dA, A(0,I), lda,
stream, &iinfo);
if ( iinfo < 0 ) {
*info = iinfo;
break;
} else if ( iinfo != 0 ) {
*info = iinfo + I * NB;
//break;
}
/* adjust pivots */
for( ii=I; ii < min(I+N,m); ii++ )
ipiv[ii] += I;
}
time_comp += timer_stop( time );
/* download the current big panel to CPU */
timer_start( time );
magmablas_zgetmatrix_transpose_mgpu(ngpu, stream, dAT, ldn_local, A(0,I), lda, dA, maxm, M, N, nb);
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
time_get += timer_stop( time );
} /* end of for */
timer_stop( time_total );
flops = FLOPS_ZGETRF( m, n ) / 1e9;
timer_printf(" memory-allocation time: %e\n", time_alloc );
timer_printf(" NB=%d nb=%d\n", (int) NB, (int) nb );
timer_printf(" memcopy and transpose %e seconds\n", time_set );
timer_printf(" total time %e seconds\n", time_total );
timer_printf(" Performance %f GFlop/s, %f seconds without htod and dtoh\n", flops / (time_comp), time_comp );
timer_printf(" Performance %f GFlop/s, %f seconds with htod\n", flops / (time_comp + time_set), time_comp + time_set );
timer_printf(" Performance %f GFlop/s, %f seconds with dtoh\n", flops / (time_comp + time_get), time_comp + time_get );
timer_printf(" Performance %f GFlop/s, %f seconds without memory-allocation\n", flops / (time_total - time_alloc), time_total - time_alloc );
for( d=0; d < ngpu0; d++ ) {
magma_setdevice(d);
magma_free( dA[d] );
magma_event_destroy( event[d][0] );
magma_event_destroy( event[d][1] );
magma_queue_destroy( stream[d][0] );
magma_queue_destroy( stream[d][1] );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
if ( *info >= 0 )
magma_zgetrf_piv(m, n, NB, A, lda, ipiv, info);
return *info;
} /* magma_zgetrf_m */
extern "C" magma_int_t
magma_zgetrf2_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n, magma_int_t nb, magma_int_t offset,
cuDoubleComplex **d_lAT, magma_int_t lddat, magma_int_t *ipiv,
cuDoubleComplex **d_lAP, cuDoubleComplex *w, magma_int_t ldw,
cudaStream_t streaml[][2], magma_int_t *info)
#endif
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2010
Purpose
=======
ZGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Use two buffer to send panels..
Arguments
=========
NUM_GPUS
(input) INTEGER
The number of GPUS to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX_16 array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -7, internal GPU memory allocation failed.
> 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define inAT(id,i,j) (d_lAT[(id)] + ((offset)+(i)*nb)*lddat + (j)*nb)
#define W(j) (w+((j)%num_gpus)*nb*ldw)
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t block_size = 32;
magma_int_t iinfo, n_local[4];
magma_int_t maxm, mindim;
magma_int_t i, ii, d, dd, rows, cols, s, ldpan[4];
magma_int_t id, i_local, i_local2, nb0, nb1;
cuDoubleComplex *d_panel[4], *panel_local[4];
//cudaStream_t streaml[4][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (num_gpus*lddat < max(1,n))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
//nb = magma_get_zgetrf_nb(m);
if( num_gpus > ceil((double)n/nb) ) {
*info = -1;
return *info;
}
{
/* Use hybrid blocked code. */
maxm = ((m + block_size-1)/block_size)*block_size;
/* some initializations */
for(i=0; i<num_gpus; i++){
magmaSetDevice(i);
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
/* workspaces */
d_panel[i] = &(d_lAP[i][nb*maxm]); /* temporary panel storage */
/* create local streams */
//magma_queue_create(&streaml[i][0]);
//magma_queue_create(&streaml[i][1]);
}
trace_init( 1, num_gpus, 2, (CUstream_st**)streaml );
/* start sending the panel to cpu */
nb0 = min(mindim, nb);
magmaSetDevice(0);
magmablasSetKernelStream(streaml[0][1]);
trace_gpu_start( 0, 1, "comm", "get" );
if( nb0 == nb )
magmablas_ztranspose( d_lAP[0], maxm, inAT(0,0,0), lddat, nb0, maxm );
else
magmablas_ztranspose2( d_lAP[0], maxm, inAT(0,0,0), lddat, nb0, maxm );
magma_zgetmatrix_async( m, nb0,
d_lAP[0], maxm,
W(0), ldw, streaml[0][1] );
trace_gpu_end( 0, 1 );
/* ------------------------------------------------------------------------------------- */
#ifdef PROFILE
magma_timestr_t start_timer, end_timer;
start_timer = get_current_time();
#endif
s = mindim / nb;
for( i=0; i<s; i++ )
{
/* Set the GPU number that holds the current panel */
id = i%num_gpus;
magmaSetDevice(id);
/* Set the local index where the current panel is */
i_local = i/num_gpus;
cols = maxm - i*nb;
rows = m - i*nb;
/* synchrnoize i-th panel from id-th gpu into work */
magma_queue_sync( streaml[id][1] );
/* i-th panel factorization */
trace_cpu_start( 0, "getrf", "getrf" );
#ifdef PANEL_FACT_MC
cntxt->nb = 12;
magma_zgetrf_mc(cntxt, &rows, &nb, W(i), &ldw, ipiv+i*nb, &iinfo);
#else
lapackf77_zgetrf( &rows, &nb, W(i), &ldw, ipiv+i*nb, &iinfo);
#endif
if ( (*info == 0) && (iinfo > 0) ) {
*info = iinfo + i*nb;
//break;
}
trace_cpu_end( 0 );
/* start sending the panel to all the gpus */
d = (i+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
magmaSetDevice(d);
trace_gpu_start( 0, 1, "comm", "set" );
magma_zsetmatrix_async( rows, nb,
W(i), ldw,
d_lAP[d], cols, streaml[d][1] );
trace_gpu_end( 0, 1 );
d = (d+1)%num_gpus;
}
/* apply the pivoting */
d = (i+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][0]);
trace_gpu_start( d, 1, "pivot", "pivot" );
if( dd == 0 )
magmablas_zpermute_long2( lddat, inAT(d,0,0), lddat, ipiv, nb, i*nb );
else
magmablas_zpermute_long3( inAT(d,0,0), lddat, ipiv, nb, i*nb );
trace_gpu_end( d, 1 );
d = (d+1)%num_gpus;
}
/* update the trailing-matrix/look-ahead */
d = (i+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
magmaSetDevice(d);
/* storage for panel */
if( d == id ) {
/* the panel belond to this gpu */
panel_local[d] = inAT(d,i,i_local);
ldpan[d] = lddat;
/* next column */
i_local2 = i_local+1;
} else {
/* the panel belong to another gpu */
panel_local[d] = &d_panel[d][(i%2)*nb*maxm];
//panel_local[d] = d_panel[d];
ldpan[d] = nb;
/* next column */
i_local2 = i_local;
if( d < id ) i_local2 ++;
}
/* the size of the next column */
if ( s > (i+1) ) {
nb0 = nb;
} else {
nb0 = n_local[d]-nb*(s/num_gpus);
if( d < s%num_gpus ) nb0 -= nb;
}
if( d == (i+1)%num_gpus) {
/* owns the next column, look-ahead the column */
nb1 = nb0;
magmablasSetKernelStream(streaml[d][1]);
/* make sure all the pivoting has been applied */
magma_queue_sync(streaml[d][0]);
trace_gpu_start( d, 1, "gemm", "gemm" );
} else {
/* update the entire trailing matrix */
nb1 = n_local[d] - i_local2*nb;
magmablasSetKernelStream(streaml[d][0]);
/* synchronization to make sure panel arrived on gpu */
magma_queue_sync(streaml[d][1]);
trace_gpu_start( d, 0, "gemm", "gemm" );
}
magmablas_ztranspose(panel_local[d], ldpan[d], d_lAP[d], cols, cols, nb);
/* gpu updating the trailing matrix */
//magmablas_ztrsm(
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb1, nb, c_one,
panel_local[d], ldpan[d],
inAT(d, i, i_local2), lddat);
//cublasZgemm
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
nb1, m-(i+1)*nb, nb,
c_neg_one, inAT(d, i, i_local2), lddat,
&(panel_local[d][nb*ldpan[d]]), ldpan[d],
c_one, inAT(d, i+1, i_local2), lddat );
if( d == (i+1)%num_gpus )
{
/* Set the local index where the current panel is */
int loff = i+1;
int i_local = (i+1)/num_gpus;
int ldda = maxm - (i+1)*nb;
int cols = m - (i+1)*nb;
nb0 = min(nb, mindim - (i+1)*nb); /* size of the diagonal block */
trace_gpu_end( d, 1 );
if( nb0 > 0 ) {
/* transpose the panel for sending it to cpu */
trace_gpu_start( d, 1, "comm", "get" );
if( i+1 < s )
magmablas_ztranspose( d_lAP[d], ldda, inAT(d,loff,i_local), lddat, nb0, ldda );
else
magmablas_ztranspose2( d_lAP[d], ldda, inAT(d,loff,i_local), lddat, nb0, ldda );
/* send the panel to cpu */
magma_zgetmatrix_async( cols, nb0,
d_lAP[d], ldda,
W(i+1), ldw, streaml[d][1] );
trace_gpu_end( d, 1 );
}
} else {
trace_gpu_end( d, 0 );
}
d = (d+1)%num_gpus;
}
/* update the remaining matrix by gpu owning the next panel */
if( (i+1) < s ) {
int i_local = (i+1)/num_gpus;
int rows = m - (i+1)*nb;
d = (i+1)%num_gpus;
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][0]);
trace_gpu_start( d, 0, "gemm", "gemm" );
//magmablas_ztrsm
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n_local[d] - (i_local+1)*nb, nb,
c_one, panel_local[d], ldpan[d],
inAT(d,i,i_local+1), lddat );
//cublasZgemm
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n_local[d]-(i_local+1)*nb, rows, nb,
c_neg_one, inAT(d,i,i_local+1), lddat,
&(panel_local[d][nb*ldpan[d]]), ldpan[d],
c_one, inAT(d,i+1, i_local+1), lddat );
trace_gpu_end( d, 0 );
}
} /* end of for i=1..s */
/* ------------------------------------------------------------------------------ */
/* Set the GPU number that holds the last panel */
id = s%num_gpus;
/* Set the local index where the last panel is */
i_local = s/num_gpus;
/* size of the last diagonal-block */
nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
cols = maxm - s*nb;
if( nb0 > 0 ) {
magmaSetDevice(id);
/* wait for the last panel on cpu */
magma_queue_sync( streaml[id][1] );
/* factor on cpu */
lapackf77_zgetrf( &rows, &nb0, W(s), &ldw, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
/* send the factor to gpus */
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
i_local2 = i_local;
if( d < id ) i_local2 ++;
if( d == id || n_local[d] > i_local2*nb ) {
magma_zsetmatrix_async( rows, nb0,
W(s), ldw,
d_lAP[d], cols, streaml[d][1] );
}
}
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][0]);
if( d == 0 )
magmablas_zpermute_long2( lddat, inAT(d,0,0), lddat, ipiv, nb0, s*nb );
else
magmablas_zpermute_long3( inAT(d,0,0), lddat, ipiv, nb0, s*nb );
}
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
magmablasSetKernelStream(streaml[d][1]);
/* wait for the pivoting to be done */
magma_queue_sync( streaml[d][0] );
i_local2 = i_local;
if( d < id ) i_local2++;
if( d == id ) {
/* the panel belond to this gpu */
panel_local[d] = inAT(d,s,i_local);
/* next column */
nb1 = n_local[d] - i_local*nb-nb0;
magmablas_ztranspose2( panel_local[d], lddat, d_lAP[d], cols, rows, nb0);
if( nb1 > 0 )
//cublasZtrsm
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb1, nb0, c_one,
panel_local[d], lddat,
inAT(d,s,i_local)+nb0, lddat);
} else if( n_local[d] > i_local2*nb ) {
/* the panel belong to another gpu */
panel_local[d] = &d_panel[d][(s%2)*nb*maxm];
//panel_local[d] = d_panel[d];
/* next column */
nb1 = n_local[d] - i_local2*nb;
magmablas_ztranspose2( panel_local[d], nb, d_lAP[d], cols, rows, nb0);
//cublasZtrsm
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb1, nb0, c_one,
panel_local[d], nb,
inAT(d,s,i_local2), lddat);
}
}
} /* if( nb0 > 0 ) */
/* clean up */
trace_finalize( "zgetrf_mgpu.svg","trace.css" );
for( d=0; d<num_gpus; d++ ) {
magmaSetDevice(d);
magma_queue_sync( streaml[d][0] );
magma_queue_sync( streaml[d][1] );
//magma_queue_destroy(streaml[d][0]);
//magma_queue_destroy(streaml[d][1]);
magmablasSetKernelStream(NULL);
}
magmaSetDevice(0);
#ifdef PROFILE
end_timer = get_current_time();
printf("\n Performance %f GFlop/s\n", (2./3.*n*n*n /1000000.) / GetTimerValue(start_timer, end_timer));
#endif
}
return *info;
/* End of MAGMA_ZGETRF2_MGPU */
}
extern "C" magma_int_t
magma_zgetrf_mgpu(magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
cuDoubleComplex **d_lA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
ZGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
=========
NUM_GPUS
(input) INTEGER
The number of GPUS to be used for the factorization.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX_16 array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
> 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define inAT(id,i,j) (d_lAT[(id)] + (i)*nb*lddat + (j)*nb)
cuDoubleComplex c_one = MAGMA_Z_ONE;
cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo, nb, n_local[MagmaMaxGPUs];
magma_int_t maxm, mindim;
magma_int_t i, j, d, rows, cols, s, lddat, lddwork;
magma_int_t id, i_local, i_local2, nb0, nb1;
cuDoubleComplex *d_lAT[MagmaMaxGPUs];
cuDoubleComplex *d_panel[MagmaMaxGPUs], *work;
cudaStream_t streaml[4][2];
/* Check arguments */
*info = 0;
if (m < 0)
*info = -2;
else if (n < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_zgetrf_nb(m);
if (nb <= 1 || nb >= n) {
/* Use CPU code. */
magma_zmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_zgetmatrix( m, n, d_lA[0], ldda, work, m );
lapackf77_zgetrf(&m, &n, work, &m, ipiv, info);
magma_zsetmatrix( m, n, work, m, d_lA[0], ldda );
magma_free_cpu(work);
} else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if( num_gpus > ceil((double)n/nb) ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) num_gpus );
*info = -1;
return *info;
}
/* allocate workspace for each GPU */
lddat = ((((((n+nb-1)/nb)/num_gpus)*nb)+31)/32)*32;
lddat = (n+nb-1)/nb; /* number of block columns */
lddat = (lddat+num_gpus-1)/num_gpus; /* number of block columns per GPU */
lddat = nb*lddat; /* number of columns per GPU */
lddat = ((lddat+31)/32)*32; /* make it a multiple of 32 */
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
/* local-n and local-ld */
n_local[i] = ((n/nb)/num_gpus)*nb;
if (i < (n/nb)%num_gpus)
n_local[i] += nb;
else if (i == (n/nb)%num_gpus)
n_local[i] += n%nb;
/* workspaces */
if (MAGMA_SUCCESS != magma_zmalloc( &d_panel[i], 3*nb*maxm )) {
for( j=0; j<=i; j++ ) {
magma_setdevice(j);
}
for( j=0; j<i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* local-matrix storage */
if (MAGMA_SUCCESS != magma_zmalloc( &d_lAT[i], lddat*maxm )) {
for( j=0; j<=i; j++ ) {
magma_setdevice(j);
magma_free( d_panel[j] );
}
for( j=0; j<i; j++ ) {
magma_setdevice(j);
magma_free( d_lAT[j] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* create the streams */
magma_queue_create( &streaml[i][0] );
magma_queue_create( &streaml[i][1] );
magmablasSetKernelStream(streaml[i][1]);
magmablas_ztranspose2( d_lAT[i], lddat, d_lA[i], ldda, m, n_local[i] );
}
for(i=0; i<num_gpus; i++){
magma_setdevice(i);
cudaStreamSynchronize(streaml[i][0]);
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
/* cpu workspace */
lddwork = maxm;
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, lddwork*nb*num_gpus )) {
for(i=0; i<num_gpus; i++ ) {
magma_setdevice(i);
magma_free( d_panel[i] );
magma_free( d_lAT[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* calling multi-gpu interface with allocated workspaces and streams */
//magma_zgetrf1_mgpu( num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
// (cudaStream_t **)streaml, info );
magma_zgetrf2_mgpu(num_gpus, m, n, nb, 0, d_lAT, lddat, ipiv, d_panel, work, maxm,
streaml, info);
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
/* save on output */
magmablas_ztranspose2( d_lA[d], ldda, d_lAT[d], lddat, n_local[d], m );
magma_device_sync();
magma_free( d_lAT[d] );
magma_free( d_panel[d] );
magma_queue_destroy( streaml[d][0] );
magma_queue_destroy( streaml[d][1] );
magmablasSetKernelStream(NULL);
} /* end of for d=1,..,num_gpus */
magma_setdevice(0);
magma_free_pinned( work );
}
return *info;
/* End of MAGMA_ZGETRF_MGPU */
}
/**
Purpose
-------
ZGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_zgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_zgetrf_gpu(
magma_int_t m, magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda, magma_int_t *ipiv,
magma_int_t *info)
{
#define dAT(i_, j_) (dAT + (i_)*nb*lddat + (j_)*nb)
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, j, rows, cols, s, lddat, ldwork;
magmaDoubleComplex *dAT, *dAP, *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_zgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_zmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_zgetmatrix( m, n, dA, ldda, work, m );
lapackf77_zgetrf(&m, &n, work, &m, ipiv, info);
magma_zsetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
if (MAGMA_SUCCESS != magma_zmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// square matrices can be done in place;
// rectangular requires copy to transpose
if ( m == n ) {
dAT = dA;
lddat = ldda;
magmablas_ztranspose_inplace( m, dAT, ldda );
}
else {
lddat = maxn; // N-by-M
if (MAGMA_SUCCESS != magma_zmalloc( &dAT, lddat*maxm )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ztranspose( m, n, dA, ldda, dAT, lddat );
}
ldwork = maxm;
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, ldwork*nb )) {
magma_free( dAP );
if ( ! (m == n))
magma_free( dAT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( j=0; j < s; j++ ) {
// download j-th panel
cols = maxm - j*nb;
magmablas_ztranspose( nb, m-j*nb, dAT(j,j), lddat, dAP, cols );
// make sure that the transpose has completed
magma_queue_sync( stream[1] );
magma_zgetmatrix_async( m-j*nb, nb, dAP, cols, work, ldwork,
stream[0]);
if ( j > 0 ) {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (j+1)*nb, nb,
c_one, dAT(j-1,j-1), lddat,
dAT(j-1,j+1), lddat );
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-j*nb, nb,
c_neg_one, dAT(j-1,j+1), lddat,
dAT(j, j-1), lddat,
c_one, dAT(j, j+1), lddat );
}
// do the cpu part
rows = m - j*nb;
magma_queue_sync( stream[0] );
lapackf77_zgetrf( &rows, &nb, work, &ldwork, ipiv+j*nb, &iinfo);
if ( *info == 0 && iinfo > 0 )
*info = iinfo + j*nb;
// upload j-th panel
magma_zsetmatrix_async( m-j*nb, nb, work, ldwork, dAP, maxm,
stream[0]);
for( i=j*nb; i < j*nb + nb; ++i ) {
ipiv[i] += j*nb;
}
magmablas_zlaswp( n, dAT, lddat, j*nb + 1, j*nb + nb, ipiv, 1 );
magma_queue_sync( stream[0] );
magmablas_ztranspose( m-j*nb, nb, dAP, maxm, dAT(j,j), lddat );
// do the small non-parallel computations (next panel update)
if ( s > (j+1) ) {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat);
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat );
}
else {
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat);
magma_zgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ztranspose( nb0, rows, dAT(s,s), lddat, dAP, maxm );
magma_zgetmatrix( rows, nb0, dAP, maxm, work, ldwork );
// do the cpu part
lapackf77_zgetrf( &rows, &nb0, work, &ldwork, ipiv+s*nb, &iinfo);
if ( *info == 0 && iinfo > 0 )
*info = iinfo + s*nb;
for( i=s*nb; i < s*nb + nb0; ++i ) {
ipiv[i] += s*nb;
}
magmablas_zlaswp( n, dAT, lddat, s*nb + 1, s*nb + nb0, ipiv, 1 );
// upload j-th panel
magma_zsetmatrix( rows, nb0, work, ldwork, dAP, maxm );
magmablas_ztranspose( rows, nb0, dAP, maxm, dAT(s,s), lddat );
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
}
// undo transpose
if ( m == n ) {
magmablas_ztranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ztranspose( n, m, dAT, lddat, dA, ldda );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_zgetrf_gpu */
/**
Purpose
-------
DORGQR generates an M-by-N DOUBLE_PRECISION matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF.
This version recomputes the T matrices on the CPU and sends them to the GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
A DOUBLE_PRECISION array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by DGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
lda INTEGER
The first dimension of the array A. LDA >= max(1,M).
@param[in]
tau DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF_GPU.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument has an illegal value
@ingroup magma_dgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_dorgqr2(magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
magma_int_t *info)
{
#define A(i,j) ( A + (i) + (j)*lda )
#define dA(i,j) (dA + (i) + (j)*ldda)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
magma_int_t nb = magma_get_dgeqrf_nb(min(m, n));
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk; //, iinfo;
magma_int_t lddwork;
double *dA, *dV, *dW, *dT, *T;
double *work;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
} else {
ki = 0;
kk = 0;
}
// Allocate GPU work space
// ldda*n for matrix dA
// ldda*nb for dV
// lddwork*nb for dW larfb workspace
ldda = ((m + 31) / 32) * 32;
lddwork = ((n + 31) / 32) * 32;
if (MAGMA_SUCCESS != magma_dmalloc( &dA, ldda*n + ldda*nb + lddwork*nb + nb*nb)) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dV = dA + ldda*n;
dW = dA + ldda*n + ldda*nb;
dT = dA + ldda*n + ldda*nb + lddwork*nb;
// Allocate CPU work space
lwork = (n+m+nb) * nb;
magma_dmalloc_cpu( &work, lwork );
T = work;
if (work == NULL) {
magma_free( dA );
magma_free_cpu( work );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
double *V = work + (n+nb)*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
/*
lapackf77_dorgqr( &m_kk, &n_kk, &k_kk,
A(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo );
*/
lapackf77_dlacpy( MagmaUpperLowerStr, &m_kk, &k_kk, A(kk,kk), &lda, V, &m_kk);
lapackf77_dlaset( MagmaUpperLowerStr, &m_kk, &n_kk, &c_zero, &c_one, A(kk, kk), &lda );
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &k_kk,
V, &m_kk, &tau[kk], work, &k_kk);
lapackf77_dlarfb( MagmaLeftStr, MagmaNoTransStr, MagmaForwardStr, MagmaColumnwiseStr,
&m_kk, &n_kk, &k_kk,
V, &m_kk, work, &k_kk, A(kk, kk), &lda, work+k_kk*k_kk, &n_kk );
if (kk > 0) {
magma_dsetmatrix( m_kk, n_kk,
A(kk, kk), lda,
dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_dlaset( MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda );
}
}
if (kk > 0) {
// Use blocked code
// stream: set Aii (V) --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min(nb, k - i);
// Send current panel to the GPU
mi = m - i;
lapackf77_dlaset( "Upper", &ib, &ib, &c_zero, &c_one, A(i, i), &lda );
magma_dsetmatrix_async( mi, ib,
A(i, i), lda,
dV, ldda, stream );
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&mi, &ib,
A(i,i), &lda, &tau[i], T, &nb);
magma_dsetmatrix_async( ib, ib,
T, nb,
dT, nb, stream );
// set panel to identity
magmablas_dlaset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda );
magmablas_dlaset( MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda );
magma_queue_sync( stream );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT, nb,
dA(i, i), ldda, dW, lddwork );
}
}
// copy result back to CPU
magma_dgetmatrix( m, n,
dA(0, 0), ldda, A(0, 0), lda);
}
magmablasSetKernelStream( NULL );
magma_queue_destroy( stream );
magma_free( dA );
magma_free_cpu( work );
return *info;
} /* magma_dorgqr */
extern "C" magma_int_t
magma_cgetrf_gpu(magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
> 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define dAT(i,j) (dAT + (i)*nb*lddat + (j)*nb)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, rows, cols, s, lddat, lddwork;
magmaFloatComplex *dAT, *dAP, *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
lddat = maxn;
lddwork = maxm;
dAT = dA;
if (MAGMA_SUCCESS != magma_cmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if ( m == n ) {
lddat = ldda;
magmablas_ctranspose_inplace( m, dAT, ldda );
}
else {
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ctranspose2( dAT, lddat, dA, ldda, m, n );
}
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
magma_free( dAP );
if ( ! (m == n))
magma_free( dAT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
cudaStream_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else
stream[1] = current_stream;
for( i=0; i<s; i++ )
{
// download i-th panel
cols = maxm - i*nb;
//magmablas_ctranspose( dAP, cols, dAT(i,i), lddat, nb, cols );
magmablas_ctranspose2( dAP, cols, dAT(i,i), lddat, nb, m-i*nb );
// make sure that that the transpose has completed
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dAP, cols, work, lddwork,
stream[0]);
if ( i>0 ){
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), lddat,
dAT(i-1,i+1), lddat );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), lddat,
dAT(i, i-1), lddat,
c_one, dAT(i, i+1), lddat );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lddwork, ipiv+i*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dAP, maxm,
stream[0]);
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
//magmablas_ctranspose(dAT(i,i), lddat, dAP, maxm, cols, nb);
magmablas_ctranspose2(dAT(i,i), lddat, dAP, maxm, m-i*nb, nb);
// do the small non-parallel computations (next panel update)
if ( s > (i+1) ) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose2( dAP, maxm, dAT(s,s), lddat, nb0, rows);
magma_cgetmatrix( rows, nb0, dAP, maxm, work, lddwork );
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lddwork, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb0, s*nb );
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dAP, maxm );
magmablas_ctranspose2( dAT(s,s), lddat, dAP, maxm, rows, nb0);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
if ( m == n ) {
magmablas_ctranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ctranspose2( dA, ldda, dAT, lddat, n, m );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* End of MAGMA_CGETRF_GPU */
void magmablas_ssymm_mgpu_com(
magma_side_t side, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
float alpha,
float *dA[], magma_int_t ldda, magma_int_t offset,
float *dB[], magma_int_t lddb,
float beta, float *dC[], magma_int_t lddc,
float *dwork[], magma_int_t dworksiz,
float *C, magma_int_t ldc,
float *work[], magma_int_t worksiz,
magma_int_t ngpu, magma_int_t nb,
magma_queue_t streams[][20], magma_int_t nstream,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10], magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
//printf("####################################################\n");
//printf(" start ssymm \n");
//printf("####################################################\n");
if ( side != MagmaLeft || uplo != MagmaLower ) {
fprintf( stderr, "%s: only Left Lower implemented\n", __func__ );
}
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nstream >= ngpu );
assert( nbevents >= ngpu*ngpu );
float c_one = MAGMA_S_ONE;
float *dwork1[MagmaMaxGPUs];
float *dwork2[MagmaMaxGPUs];
magma_int_t maxgsize = n*m;
magma_int_t lddwork = lddc;
magma_int_t ldwork = m;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev]; // size of dwork1 is n*lddwork
dwork2[dev] = dwork[dev]+n*lddwork; // size of dwork2 is maxgsize*ngpu
}
assert( dworksiz >= (n*lddwork+maxgsize*ngpu) );
assert( worksiz >= (n*ldwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev, devperm, myblk, mycolsize, myblkoffst;
magma_int_t gmaster;
magma_int_t masterdev, lcdev, lccolsize, myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m, nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
cudaMemset(dwork(dev,0,0), 0, (lddwork)*(n)*sizeof(float) );
// put all dC on all dev to 0 except the one which
// hold i==0 because this one has to multiply by beta.
if(dev!=stdev){
cudaMemset(dC(dev,0,0), 0, (lddc)*(n)*sizeof(float) );
}
}
magma_int_t newoffset = offset;
// 1. symmetrize
if(blockoffset>0){
newoffset = offset+fstblksiz; // newoffset is adjusted over nb
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > stdev?1:0);
//printf("STDEV %d voici offset %d remm %d myblockoffset %d siz %d \n", stdev, offset, remm, myblkoffst, fstblksiz);
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
magmablas_ssymmetrize_tiles( MagmaLower, fstblksiz, dA(stdev, offset, myblkoffst*nb+blockoffset), ldda, 1, ngpu*nb, nb );
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t newstdev = (newoffset/nb)%ngpu;
magma_int_t nbblk = remm/nb; // number of block of size nb. if m%nb>0 then a last block exist and is of size ib=m%nb
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-newstdev+ngpu)%ngpu) ? 1:0 );
magma_int_t devperm = (dev-newstdev+ngpu)%ngpu;
magma_int_t nbblkoffst = newoffset/nb;
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
//printf("dev %d devperm %d newoffset %d rowoff %d coloff %d myblk %d \n", dev, devperm, newoffset, newoffset+devperm*nb, myblkoffst*nb, myblk);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magmablas_ssymmetrize_tiles( MagmaLower, nb, dA(dev, newoffset+devperm*nb, myblkoffst*nb), ldda, myblk, ngpu*nb, nb );
if(remm%nb>0){
magma_int_t nblstblks = (nbblk+1)%ngpu;
magma_int_t devlstblk = (nblstblks-1+ngpu)%ngpu;
//printf("==> siz %d devperm %d, devlstblk %d, newoffset+nbblk*nb %d, myblkoffst*nb+ myblk*nb %d\n", remm % nb, devperm, devlstblk, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb);
if(devperm==devlstblk)
magmablas_ssymmetrize( MagmaLower, remm % nb, dA(dev, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb), ldda ); // last partial tile
}
}
/*
magma_int_t siz = m+offset;
float *R;
magma_smalloc_cpu( &R, siz*siz );
// collecte back A
magmablas_sgetmatrix_1D_bcyclic( siz, siz, dA, ldda, R, siz, ngpu, nb );
magma_setdevice( 0 );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
//magma_sgetmatrix( siz, siz, dA[0], ldda, R, siz );
FILE *trace_file;
trace_file = fopen("AJETE/Aafter", "w");
for (int j = 0; j < siz ; j++)
for (int i = 0; i < siz ; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, R[j*siz+i]);
fclose(trace_file);
return;
*/
// ROW GEMM transpose a row and make a gemm with a block
// if only 1 GPU used the ROW GEMM is integrated with the
// COL GEMM (better accuracy observed) and better perf
if(ngpu>1){
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
//magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nbblk = magma_ceildiv(i, nb);
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-stdev+ngpu)%ngpu) ? 1:0 );
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
magma_int_t myrowsize = myblk * nb;
magma_int_t coloffset = myblkoffst*nb;
if(dev==stdev) {
myrowsize = myrowsize -blockoffset;
coloffset = myblkoffst*nb+blockoffset;
}
//printf("ROW GEMM: voici i %d ib %d ioff %d nbblkoffst %d stdev %d dev %d myblk %d myblkoffset %d coloffset %d rowsize %d\n", i, ib, ioff, nbblkoffst, stdev, dev, myblk, myblkoffst, coloffset, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, myrowsize, n, ib,
alpha, dA(dev,ioff,coloffset), ldda,
dB(dev,i,0), lddb,
c_one, dwork(dev,0,0), lddwork );
}
}
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_event_record(redevents[dev][1], streams[dev][1]);
}
}
// COL GEMM
// blockoffset is offset within first block; for subsequent blocks it is 0
if(blockoffset>0){
magma_int_t ib = min( nb-blockoffset, m ); // block size
magma_int_t iblock = (offset / nb) / ngpu; // local block id
magma_int_t di = iblock*nb+blockoffset; // local index in parent matrix
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
//printf("DEV %d COL GEMM first ioff %d di %d m %d n %d ib %d \n", stdev, offset, di, m, n, ib);
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m, n, ib,
alpha, dA(stdev,offset,di), ldda,
dB(stdev,0,0), lddb,
beta, dC(stdev,0,0), lddc );
}
// COL GEMM
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
magma_int_t iblock = (ioff / nb) / ngpu; // local block id
magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t di = iblock*nb; // local index in parent matrix
//printf("DEV %d COL GEMM i %d ioff %d di %d m-i %d n %d ib %d \n", dev, i, ioff, di, m-i, n, ib);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
if(i==0){
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
beta, dC(dev,i,0), lddc );
}else{
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,i,0), lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
// if only 1 GPU is used, do the ROW GEMM
if(ngpu==1){
// NOTE THAT because the COL gemm write dC below the diagonal (i)
// and the ROW GEMM write dC from 0 to diag-1, so they could
// run in parallel on different streams.
//
// NO NO NO because
// it might happen that col finished i and strated i+1 while row still at i
// magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, i, n, ib,
alpha, dA(dev,ioff,offset), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,0,0), lddc );
}
}
if(ngpu>1){
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t nbblkrow = nbblk-1;
magma_int_t devperm = (dev-stdev+ngpu)%ngpu;
magma_int_t myblk = (nbblkrow/ngpu) + (nbblkrow%ngpu > devperm ? 1:0 );
magma_int_t myrowsize = myblk * nb;
if(dev==stdev) {
myrowsize = myrowsize - blockoffset;
}
//printf("blockoffset %d nbblkrow %d devperm %d DEV %d RECEIVING myblk %d myrowsize %d\n", blockoffset, nbblkrow, devperm, dev, myblk, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][1]);
//magma_queue_sync( streams[ dev ][ 1 ] );
// for each dev add the computed ROW block each on its placment with dC
for( magma_int_t blki = 0; blki < myblk; ++blki){
magma_int_t gbblki = (blki*ngpu + devperm)*nb - blockoffset;
magma_int_t lcblki = blki*nb;
magma_int_t ib = nb;// min(nb, m-gbblki);
if(dev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
magmablas_sgeadd(ib, n, c_one,
&dwork[dev][lcblki], lddwork,
&dC[dev][gbblki] , lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
}
}
}
// ===========================================================
// COMMUNICATION ALL_REDUCE_SUM
// ===========================================================
if(ngpu==1){
return;
}
// INITIALIZE COMM
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize, m);
if(mycolsize>0){
gpuisactive[dev] = mycolsize;
if(masterdev==-1) {
masterdev = dev;
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//*******************************
// each GPU send its result
// to its master. The master make
// the addition and then send to
// to the masters of other real
// and receive from the masters of
// other real make the addition
// and broadcast locally the final
// result.
//*******************************
//printf("=======================================================================\n");
//printf(" sending to my master \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU. if I am not the master, then send my result to my master.
// store result on dwork[masterdev][dev*maxgsize]
if(dev!=masterdev){
magma_setdevice( dev );
//printf(" GPU %d sending to my master %d\n", dev, masterdev);
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][0]);
// sending to the master of my real
magma_scopymatrix_async(
m, n,
&dC[dev][0], lddc,
&dwork2[masterdev][maxgsize*dev], m, streams[dev][0] );
magma_event_record(redevents[dev][masterdev], streams[dev][0]);
} // end I am not the masterdev
}// end if mycolsize>0
}// for idev
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//printf("=======================================================================\n");
//printf(" each master do addition of local result and broadcast to other masters \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][0]);
// ========================================
// local addition
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" master %d receiving from %d and adding \n", masterdev, lcdev);
// this is an active GPU of my real.
// wait I received what he send it to me and then do addition.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[lcdev][masterdev]);
magmablas_sgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*lcdev], m,
&dC[masterdev][0] , lddc );
}
}// for l=1:myngpu
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
//
// ========================================
// send to other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//Master has to wait until finish the local addition then send using gmaster stream.
//use stream 0 to make it sequential or stream gmaster to make it parallel.
//Now both re the same.
//printf(" master %d from cmplx %d sending to other master %d on cmplx %d \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ gmaster ], redevents[masterdev][masterdev]);
magma_scopymatrix_async(
m, n,
&dC[masterdev][0], lddc,
&dwork2[gmaster][maxgsize*masterdev], m, streams[masterdev][gmaster] );
magma_event_record(redevents[masterdev][gmaster], streams[masterdev][gmaster]);
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][gmaster]);
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// ========================================
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//printf("=======================================================================\n");
//printf(" each master wait receiving other masters results, do the addition and broadcast locally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// master has to wait until finishing all the send to other masters.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
// ========================================
// addition of results from other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//Master has to wait until receiving from gmaster, then do addition using stream 0
//printf(" master %d from cmplx %d receiving from other master %d on cmplx %d and adding \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[gmaster][masterdev]);
magmablas_sgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*gmaster], m,
&dC[masterdev][0] , lddc );
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
// ========================================
// local broadcast of final results
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
// this is an active GPU of my real.
// wait the previous addition is done maening stream 0 is finished and broadcast sequentially for now.
// to make it parallel put stream lcdev instead of stream 0
//printf(" master %d broadcasting local to %d \n", masterdev, lcdev);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
magma_scopymatrix_async(
m, n,
&dC[masterdev][0], lddc,
&dC[lcdev][0], lddc, streams[masterdev][0] );
magma_event_record(redevents[masterdev][lcdev], streams[masterdev][0]);
}
}// for l=1:myngpu
// ========================================
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if(lccolsize>0){
magma_setdevice( lcdev );
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[lcdev][0]);
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[masterdev][lcdev]);
}
}// for l=1:myngpu
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
//printf("****************************************************\n");
//printf(" finish ssymm \n");
//printf("****************************************************\n");
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}
extern "C" void magma_zbulge_applyQ(
magma_int_t WANTZ, magma_side_t SIDE, magma_int_t NE, magma_int_t N, magma_int_t NB,
magma_int_t Vblksiz, magmaDoubleComplex *E, magma_int_t LDE,
magmaDoubleComplex *V, magmaDoubleComplex *TAU, magmaDoubleComplex *T,
magma_int_t *INFO, magmaDoubleComplex *dV, magmaDoubleComplex *dT,
magmaDoubleComplex *dE, magma_int_t copytype )
{
//%===========================
//% local variables
//%===========================
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t LDT, LDV, firstcolj;
magma_int_t bg, nbGblk, rownbm, k, m, n;
magma_int_t st, ed, fst, vlen, vnb, colj, len;
magma_int_t blkid, vpos, taupos, tpos;
//magmaDoubleComplex *WORK;
magma_int_t LWORK;
magma_int_t cur_blksiz, avai_blksiz, ncolinvolvd;
magma_int_t nbgr, colst, coled, versionL, versionR;
magma_int_t blkcnt=-1;
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
*INFO=0;
versionL = 113;
versionR = 92;
LDT = Vblksiz;
LDV = NB+Vblksiz-1;
//blklen = LDV*Vblksiz;
nbGblk = plasma_ceildiv((N-1), Vblksiz);
//magma_zmalloc_cpu( &WORK, LWORK );
/* find the size of the matrix T V*/
findVTsiz(N, NB, Vblksiz, &blkcnt, &LDV);
/* Copy E & V & T to the GPU in dE and dV and dT
* depending on copytype:
* 1: mean copy only V
* 2: mean copy V and T
* 3: mean copy V, T and E
* */
if (copytype > 0) magma_zsetmatrix( LDV, blkcnt*Vblksiz, V, LDV, dV, LDV );
if (copytype > 1) magma_zsetmatrix( LDT, blkcnt*Vblksiz, T, LDT, dT, LDT );
if (copytype > 2) magma_zsetmatrix( N, NE, E, N, dE, N );
magmaDoubleComplex *dwork;
//ldwork = NE;
LWORK = 2*N*max(Vblksiz, 64);
if (MAGMA_SUCCESS != magma_zmalloc( &dwork, LWORK )) {
printf ("!!!! magma_zbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
/* WANTZ = 1 meaning E is IDENTITY so form Q using optimized update.
* So we use the reverse order from small q to large one,
* so from q_n to q_1 so Left update to Identity.
* Use versionL 113 because in 114 we need to update the whole matrix and not in icreasing order.
* WANTZ = 2 meaning E is a full matrix and need to be updated from Left or Right so use normal update
* */
if (WANTZ == 1) {
versionL=113;
SIDE = MagmaLeft;
//set the matrix to Identity here to avoid copying it from the CPU
magmablas_zlaset( MagmaFull, N, N, c_zero, c_one, dE, N );
}
printf(" APPLY Q_v115 GPU with N %d NB %d Vblksiz %d SIDE %c versionL %d versionR %d WANTZ %d \n",
(int) N, (int) NB, (int) Vblksiz, SIDE, (int) versionL, (int) versionR, (int) WANTZ);
#if defined(USESTREAM)
magma_int_t N2=N/2;
magma_int_t N1=N-N2;
printf("using stream\n");
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
#endif
if (SIDE == MagmaLeft) {
if (versionL == 113) {
for (bg = nbGblk; bg > 0; bg--) {
firstcolj = (bg-1)*Vblksiz + 1;
if (bg == nbGblk)
rownbm = plasma_ceildiv((N-(firstcolj)), NB); // last blk has size=1 used for complex to handle A(N,N-1)
else
rownbm = plasma_ceildiv((N-(firstcolj+1)), NB);
for (m = rownbm; m > 0; m--) {
vlen = 0;
vnb = 0;
colj = (bg-1)*Vblksiz; // for k=0; I compute the fst and then can remove it from the loop
fst = (rownbm -m)*NB+colj +1;
for (k=0; k < Vblksiz; k++) {
colj = (bg-1)*Vblksiz + k;
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=k+1;
}
colst = (bg-1)*Vblksiz;
findVTpos(N, NB, Vblksiz, colst, fst, &vpos, &taupos, &tpos, &blkid);
printf("voici bg %d m %d vlen %d vnb %d fcolj %d vpos %d taupos %d \n", (int) bg, (int) m, (int) vlen, (int) vnb, (int) colst+1, (int) vpos+1, (int) taupos+1);
if ((vlen > 0) && (vnb > 0)) {
if (WANTZ == 1) {
len = N-colst;
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, len, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,colst), LDE, dwork, len);
} else {
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, NE, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,0), LDE, dwork, NE);
}
}
}
}
} else if (versionL == 114) {
rownbm = plasma_ceildiv((N-1), NB);
for (m = rownbm; m > 0; m--) {
ncolinvolvd = min(N-1, m*NB);
avai_blksiz=min(Vblksiz, ncolinvolvd);
nbgr = plasma_ceildiv(ncolinvolvd, avai_blksiz);
for (n = nbgr; n > 0; n--) {
vlen = 0;
vnb = 0;
cur_blksiz = min(ncolinvolvd-(n-1)*avai_blksiz, avai_blksiz);
colst = (n-1)*avai_blksiz;
coled = colst + cur_blksiz -1;
fst = (rownbm -m)*NB+colst +1;
for (colj=colst; colj <= coled; colj++) {
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=vnb+1;
}
findVTpos(N, NB, Vblksiz, colst, fst, &vpos, &taupos, &tpos, &blkid);
//printf("voici bg %d m %d vlen %d vnb %d fcolj %d vpos %d taupos %d \n", bg, m, vlen, vnb, colst+1, vpos+1, taupos+1);
if ((vlen > 0) && (vnb > 0)) {
#if defined(USESTREAM)
magmablasSetKernelStream(stream[0]);
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, N1, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,0), LDE, dwork, N1);
magmablasSetKernelStream(stream[1]);
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, N2, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,N1), LDE, &dwork[N1*Vblksiz], N2);
#else
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, vlen, NE, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(fst,0), LDE, dwork, NE);
#endif
}
}
}
}
} else if (SIDE == MagmaRight) {
if (versionR == 91) {
for (bg =1; bg <= nbGblk; bg++) {
firstcolj = (bg-1)*Vblksiz + 1;
rownbm = plasma_ceildiv((N-(firstcolj+1)), NB);
if (bg == nbGblk) rownbm = plasma_ceildiv((N-(firstcolj)), NB); // last blk has size=1 used for complex to handle A(N,N-1)
for (m = 1; m <= rownbm; m++) {
vlen = 0;
vnb = 0;
// for k=0; I compute the fst and then can remove it from the loop
colj = (bg-1)*Vblksiz;
fst = (rownbm -m)*NB+colj +1;
for (k=0; k < Vblksiz; k++) {
colj = (bg-1)*Vblksiz + k;
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=k+1;
}
colj = (bg-1)*Vblksiz;
findVTpos(N, NB, Vblksiz, colj, fst, &vpos, &taupos, &tpos, &blkid);
//printf("voici bg %d m %d vlen %d vnb %d fcolj %d vpos %d taupos %d \n", bg, m, vlen, vnb, colj, vpos, taupos);
if ((vlen > 0) && (vnb > 0)) {
#if defined(USESTREAM)
magmablasSetKernelStream(stream[0]);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N1, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, N1);
magmablasSetKernelStream(stream[1]);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N2, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(N1, fst), LDE, &dwork[N1*Vblksiz], N2);
#else
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, NE);
#endif
}
}
}
} else if (versionR == 92) {
rownbm = plasma_ceildiv((N-1), NB);
for (m = 1; m <= rownbm; m++) {
ncolinvolvd = min(N-1, m*NB);
avai_blksiz=min(Vblksiz, ncolinvolvd);
nbgr = plasma_ceildiv(ncolinvolvd, avai_blksiz);
for (n = 1; n <= nbgr; n++) {
vlen = 0;
vnb = 0;
cur_blksiz = min(ncolinvolvd-(n-1)*avai_blksiz, avai_blksiz);
colst = (n-1)*avai_blksiz;
coled = colst + cur_blksiz -1;
fst = (rownbm -m)*NB+colst +1;
for (colj=colst; colj <= coled; colj++) {
st = (rownbm -m)*NB+colj +1;
ed = min(st+NB-1, N-1);
if (st > ed) break;
if ((st == ed) && (colj != N-2)) break;
vlen=ed-fst+1;
vnb=vnb+1;
}
findVTpos(N, NB, Vblksiz, colst, fst, &vpos, &taupos, &tpos, &blkid);
if ((vlen > 0) && (vnb > 0)) {
#if defined(USESTREAM)
magmablasSetKernelStream(stream[0]);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N1, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, N1);
magmablasSetKernelStream(stream[1]);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, N2, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(N1, fst), LDE, &dwork[N1*Vblksiz], N2);
#else
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, vlen, vnb, dV(vpos), LDV, dT(tpos), LDT, dE(0, fst), LDE, dwork, NE);
#endif
}
}
}
}
} else {
printf("ERROR SIDE %d\n", SIDE);
}
#if defined(USESTREAM)
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
#endif
magmablasSetKernelStream( orig_stream );
}
/**
Purpose
-------
DORMQR overwrites the general real M-by-N matrix C with
@verbatim
SIDE = MagmaLeft SIDE = MagmaRight
TRANS = MagmaNoTrans: Q * C C * Q
TRANS = MagmaTrans: Q**H * C C * Q**H
@endverbatim
where Q is a real unitary matrix defined as the product of k
elementary reflectors
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF. Q is of order M if SIDE = MagmaLeft and of order N
if SIDE = MagmaRight.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
side magma_side_t
- = MagmaLeft: apply Q or Q**H from the Left;
- = MagmaRight: apply Q or Q**H from the Right.
@param[in]
trans magma_trans_t
- = MagmaNoTrans: No transpose, apply Q;
- = MagmaTrans: Conjugate transpose, apply Q**H.
@param[in]
m INTEGER
The number of rows of the matrix C. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix C. N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines
the matrix Q.
If SIDE = MagmaLeft, M >= K >= 0;
if SIDE = MagmaRight, N >= K >= 0.
@param[in]
A DOUBLE_PRECISION array, dimension (LDA,K)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,k, as returned by
DGEQRF in the first k columns of its array argument A.
@param[in]
lda INTEGER
The leading dimension of the array A.
If SIDE = MagmaLeft, LDA >= max(1,M);
if SIDE = MagmaRight, LDA >= max(1,N).
@param[in]
tau DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF.
@param[in,out]
C DOUBLE_PRECISION array, dimension (LDC,N)
On entry, the M-by-N matrix C.
On exit, C is overwritten by Q*C or Q**H*C or C*Q**H or C*Q.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDC >= max(1,M).
@param[out]
work (workspace) DOUBLE_PRECISION 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.
If SIDE = MagmaLeft, LWORK >= max(1,N);
if SIDE = MagmaRight, LWORK >= max(1,M).
For optimum performance LWORK >= N*NB if SIDE = MagmaLeft, and
LWORK >= M*NB if SIDE = MagmaRight, 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
@ingroup magma_dgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_dormqr_m(
magma_int_t ngpu,
magma_side_t side, magma_trans_t trans,
magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
double *C, magma_int_t ldc,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define C(i, j) (C + (j)*ldc + (i))
#define dC(gpui, i, j) (dw[gpui] + (j)*lddc + (i))
#define dA_c(gpui, ind, i, j) (dw[gpui] + maxnlocal*lddc + (ind)*lddar*lddac + (i) + (j)*lddac)
#define dA_r(gpui, ind, i, j) (dw[gpui] + maxnlocal*lddc + (ind)*lddar*lddac + (i) + (j)*lddar)
#define dT(gpui, ind) (dw[gpui] + maxnlocal*lddc + 2*lddac*lddar + (ind)*((nb+1)*nb))
#define dwork(gpui, ind) (dw[gpui] + maxnlocal*lddc + 2*lddac*lddar + 2*((nb+1)*nb) + (ind)*(lddwork*nb))
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
const char* side_ = lapack_side_const( side );
const char* trans_ = lapack_trans_const( trans );
// TODO fix memory leak (alloc after argument checks)
magma_int_t nb = 128;
double *T;
magma_dmalloc_pinned(&T, nb*nb);
//printf("calling dormqr_m with nb=%d\n", (int) nb);
double* dw[MagmaMaxGPUs];
magma_queue_t stream [MagmaMaxGPUs][2];
magma_event_t event [MagmaMaxGPUs][2];
magma_int_t ind_c;
magma_device_t igpu;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
*info = 0;
magma_int_t left = (side == MagmaLeft);
magma_int_t notran = (trans == MagmaNoTrans);
magma_int_t lquery = (lwork == -1);
/* NQ is the order of Q and NW is the minimum dimension of WORK */
magma_int_t nq, nw;
if (left) {
nq = m;
nw = n;
} else {
nq = n;
nw = m;
}
if (! left && side != MagmaRight) {
*info = -1;
} else if (! notran && trans != MagmaTrans) {
*info = -2;
} else if (m < 0) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (k < 0 || k > nq) {
*info = -5;
} else if (lda < max(1,nq)) {
*info = -7;
} else if (ldc < max(1,m)) {
*info = -10;
} else if (lwork < max(1,nw) && ! lquery) {
*info = -12;
}
magma_int_t lwkopt = max(1,nw) * nb;
if (*info == 0) {
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0 || k == 0) {
work[0] = c_one;
return *info;
}
if (nb >= k) {
/* Use CPU code */
lapackf77_dormqr(side_, trans_, &m, &n, &k, A, &lda, tau,
C, &ldc, work, &lwork, info);
return *info;
}
magma_int_t lddc = (m+63)/64*64;
magma_int_t lddac = nq;
magma_int_t lddar = nb;
magma_int_t lddwork = nw;
magma_int_t nlocal[ MagmaMaxGPUs ] = { 0 };
magma_int_t nb_l=256;
magma_int_t nbl = (n-1)/nb_l+1; // number of blocks
magma_int_t maxnlocal = (nbl+ngpu-1)/ngpu*nb_l;
ngpu = min(ngpu, (n+nb_l-1)/nb_l); // Don't use GPU that will not have data.
magma_int_t ldw = maxnlocal*lddc // dC
+ 2*lddac*lddar // 2*dA
+ 2*(nb + 1 + lddwork)*nb; // 2*(dT and dwork)
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
if (MAGMA_SUCCESS != magma_dmalloc( &dw[igpu], ldw )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(*info) );
return *info;
}
magma_queue_create( &stream[igpu][0] );
magma_queue_create( &stream[igpu][1] );
magma_event_create( &event[igpu][0] );
magma_event_create( &event[igpu][1] );
}
/* Use hybrid CPU-MGPU code */
if (left) {
//copy C to mgpus
for (magma_int_t i = 0; i < nbl; ++i) {
magma_int_t igpu = i%ngpu;
magma_setdevice(igpu);
magma_int_t kb = min(nb_l, n-i*nb_l);
magma_dsetmatrix_async( m, kb,
C(0, i*nb_l), ldc,
dC(igpu, 0, i/ngpu*nb_l), lddc, stream[igpu][0] );
nlocal[igpu] += kb;
}
magma_int_t i1, i2, i3;
if ( !notran ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
ind_c = 0;
for (magma_int_t i = i1; (i3 < 0 ? i >= i2 : i < i2); i += i3) {
// start the copy of A panel
magma_int_t kb = min(nb, k - i);
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_event_sync(event[igpu][ind_c]); // check if the new data can be copied
magma_dsetmatrix_async(nq-i, kb,
A(i, i), lda,
dA_c(igpu, ind_c, i, 0), lddac, stream[igpu][0] );
// set upper triangular part of dA to identity
magmablas_dlaset_band_q( MagmaUpper, kb, kb, kb, c_zero, c_one, dA_c(igpu, ind_c, i, 0), lddac, stream[igpu][0] );
}
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
magma_int_t nqi = nq - i;
lapackf77_dlarft("F", "C", &nqi, &kb, A(i, i), &lda,
&tau[i], T, &kb);
/* H or H' is applied to C(1:m,i:n) */
/* Apply H or H'; First copy T to the GPU */
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_dsetmatrix_async(kb, kb,
T, kb,
dT(igpu, ind_c), kb, stream[igpu][0] );
}
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_queue_sync( stream[igpu][0] ); // check if the data was copied
magmablasSetKernelStream(stream[igpu][1]);
magma_dlarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
m-i, nlocal[igpu], kb,
dA_c(igpu, ind_c, i, 0), lddac, dT(igpu, ind_c), kb,
dC(igpu, i, 0), lddc,
dwork(igpu, ind_c), lddwork);
magma_event_record(event[igpu][ind_c], stream[igpu][1] );
}
ind_c = (ind_c+1)%2;
}
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_queue_sync( stream[igpu][1] );
}
//copy C from mgpus
for (magma_int_t i = 0; i < nbl; ++i) {
magma_int_t igpu = i%ngpu;
magma_setdevice(igpu);
magma_int_t kb = min(nb_l, n-i*nb_l);
magma_dgetmatrix( m, kb,
dC(igpu, 0, i/ngpu*nb_l), lddc,
C(0, i*nb_l), ldc );
// magma_dgetmatrix_async( m, kb,
// dC(igpu, 0, i/ngpu*nb_l), lddc,
// C(0, i*nb_l), ldc, stream[igpu][0] );
}
} else {
// TODO fix memory leak T, dw, event, stream
fprintf(stderr, "The case (side == right) is not implemented\n");
*info = MAGMA_ERR_NOT_IMPLEMENTED;
magma_xerbla( __func__, -(*info) );
return *info;
/*
if ( notran ) {
i1 = 0;
i2 = k;
i3 = nb;
} else {
i1 = (k - 1) / nb * nb;
i2 = 0;
i3 = -nb;
}
mi = m;
ic = 0;
for (i = i1; (i3 < 0 ? i >= i2 : i < i2); i += i3) {
ib = min(nb, k - i);
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
i__4 = nq - i;
lapackf77_dlarft("F", "C", &i__4, &ib, A(i, i), &lda,
&tau[i], T, &ib);
// 1) copy the panel from A to the GPU, and
// 2) set upper triangular part of dA to identity
magma_dsetmatrix( i__4, ib, A(i, i), lda, dA(i, 0), ldda );
magmablas_dlaset_band( MagmaUpper, ib, ib, ib, c_zero, c_one, dA(i, 0), ldda );
// H or H' is applied to C(1:m,i:n)
ni = n - i;
jc = i;
// Apply H or H'; First copy T to the GPU
magma_dsetmatrix( ib, ib, T, ib, dT, ib );
magma_dlarfb_gpu( side, trans, MagmaForward, MagmaColumnwise,
mi, ni, ib,
dA(i, 0), ldda, dT, ib,
dC(ic, jc), lddc,
dwork, lddwork);
}
*/
}
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
for (igpu = 0; igpu < ngpu; ++igpu) {
magma_setdevice(igpu);
magma_event_destroy( event[igpu][0] );
magma_event_destroy( event[igpu][1] );
magma_queue_destroy( stream[igpu][0] );
magma_queue_destroy( stream[igpu][1] );
magma_free( dw[igpu] );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_dormqr */
extern "C" magma_int_t
magma_ziterref(
magma_z_sparse_matrix A, magma_z_vector b, magma_z_vector *x,
magma_z_solver_par *solver_par, magma_z_preconditioner *precond_par,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
// prepare solver feedback
solver_par->solver = Magma_ITERREF;
solver_par->numiter = 0;
solver_par->info = MAGMA_SUCCESS;
double residual;
magma_zresidual( A, b, *x, &residual, queue );
solver_par->init_res = residual;
// some useful variables
magmaDoubleComplex c_zero = MAGMA_Z_ZERO, c_one = MAGMA_Z_ONE,
c_mone = MAGMA_Z_NEG_ONE;
magma_int_t dofs = A.num_rows;
// workspace
magma_z_vector r,z;
magma_z_vinit( &r, Magma_DEV, dofs, c_zero, queue );
magma_z_vinit( &z, Magma_DEV, dofs, c_zero, queue );
// solver variables
double nom, nom0, r0;
// solver setup
magma_zscal( dofs, c_zero, x->dval, 1) ; // x = 0
magma_zcopy( dofs, b.dval, 1, r.dval, 1 ); // r = b
nom0 = magma_dznrm2(dofs, r.dval, 1); // nom0 = || r ||
nom = nom0 * nom0;
solver_par->init_res = nom0;
if ( (r0 = nom * solver_par->epsilon) < ATOLERANCE )
r0 = ATOLERANCE;
if ( nom < r0 ) {
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
//Chronometry
real_Double_t tempo1, tempo2;
tempo1 = magma_sync_wtime( queue );
if ( solver_par->verbose > 0 ) {
solver_par->res_vec[0] = nom0;
solver_par->timing[0] = 0.0;
}
// start iteration
for( solver_par->numiter= 1; solver_par->numiter<solver_par->maxiter;
solver_par->numiter++ ) {
magma_zscal( dofs, MAGMA_Z_MAKE(1./nom, 0.), r.dval, 1) ; // scale it
magma_z_precond( A, r, &z, precond_par, queue ); // inner solver: A * z = r
magma_zscal( dofs, MAGMA_Z_MAKE(nom, 0.), z.dval, 1) ; // scale it
magma_zaxpy(dofs, c_one, z.dval, 1, x->dval, 1); // x = x + z
magma_z_spmv( c_mone, A, *x, c_zero, r, queue ); // r = - A x
magma_zaxpy(dofs, c_one, b.dval, 1, r.dval, 1); // r = r + b
nom = magma_dznrm2(dofs, r.dval, 1); // nom = || r ||
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) nom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
if ( nom < r0 ) {
break;
}
}
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
magma_zresidual( A, b, *x, &residual, queue );
solver_par->final_res = residual;
solver_par->iter_res = nom;
if ( solver_par->numiter < solver_par->maxiter) {
solver_par->info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res ) {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) nom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_SLOW_CONVERGENCE;
}
else {
if ( solver_par->verbose > 0 ) {
if ( (solver_par->numiter)%solver_par->verbose==0 ) {
solver_par->res_vec[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) nom;
solver_par->timing[(solver_par->numiter)/solver_par->verbose]
= (real_Double_t) tempo2-tempo1;
}
}
solver_par->info = MAGMA_DIVERGENCE;
}
magma_z_vfree(&r, queue );
magma_z_vfree(&z, queue );
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
} /* magma_ziterref */
/**
Purpose
-------
SGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_sgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_sgetrf(
magma_int_t m, magma_int_t n, float *A, magma_int_t lda,
magma_int_t *ipiv,
magma_int_t *info)
{
#define dAT(i_, j_) (dAT + (i_)*nb*ldda + (j_)*nb)
float *dAT, *dA, *da, *work;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t iinfo, nb;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
nb = magma_get_sgetrf_nb(m);
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_sgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, maxdim;
magma_int_t i, j, rows, cols, s = min(m, n)/nb;
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
/* set number of GPUs */
magma_int_t ngpu = magma_num_gpus();
if ( ngpu > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_sgetrf_m(ngpu, m, n, A, lda, ipiv, info);
return *info;
}
/* explicitly checking the memory requirement */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(float);
int h = 1+(2+ngpu), ngpu2 = ngpu;
int NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
if ( ngpu > ceil((float)NB/nb) ) {
ngpu2 = (int)ceil((float)NB/nb);
h = 1+(2+ngpu2);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
}
if ( ngpu2*NB < n ) {
/* require too much memory, so call non-GPU-resident version */
magma_sgetrf_m(ngpu, m, n, A, lda, ipiv, info);
return *info;
}
ldda = maxn;
work = A;
if (maxdim*maxdim < 2*maxm*maxn) {
// if close to square, allocate square matrix and transpose in-place
if (MAGMA_SUCCESS != magma_smalloc( &dA, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_sgetrf_m(ngpu, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
ldda = maxdim;
magma_ssetmatrix( m, n, A, lda, da, ldda );
dAT = da;
magmablas_stranspose_inplace( ldda, dAT, ldda );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
if (MAGMA_SUCCESS != magma_smalloc( &dA, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_sgetrf_m(ngpu, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
magma_ssetmatrix( m, n, A, lda, da, maxm );
if (MAGMA_SUCCESS != magma_smalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dA );
magma_sgetrf_m(ngpu, m, n, A, lda, ipiv, info);
return *info;
}
magmablas_stranspose( m, n, da, maxm, dAT, ldda );
}
lapackf77_sgetrf( &m, &nb, work, &lda, ipiv, &iinfo);
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( j = 0; j < s; j++ ) {
// download j-th panel
cols = maxm - j*nb;
if (j > 0) {
magmablas_stranspose( nb, cols, dAT(j,j), ldda, dA, cols );
// make sure that gpu queue is empty
magma_device_sync();
magma_sgetmatrix_async( m-j*nb, nb, dA, cols, work, lda,
stream[0]);
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (j+1)*nb, nb,
c_one, dAT(j-1,j-1), ldda,
dAT(j-1,j+1), ldda );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-j*nb, nb,
c_neg_one, dAT(j-1,j+1), ldda,
dAT(j, j-1), ldda,
c_one, dAT(j, j+1), ldda );
// do the cpu part
rows = m - j*nb;
magma_queue_sync( stream[0] );
lapackf77_sgetrf( &rows, &nb, work, &lda, ipiv+j*nb, &iinfo);
}
if (*info == 0 && iinfo > 0)
*info = iinfo + j*nb;
// upload j-th panel
magma_ssetmatrix_async( m-j*nb, nb, work, lda, dA, cols,
stream[0]);
for( i=j*nb; i < j*nb + nb; ++i ) {
ipiv[i] += j*nb;
}
magmablas_slaswp( n, dAT, ldda, j*nb + 1, j*nb + nb, ipiv, 1 );
magma_queue_sync( stream[0] );
magmablas_stranspose( cols, nb, dA, cols, dAT(j,j), ldda );
// do the small non-parallel computations (next panel update)
if (s > (j+1)) {
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(j, j ), ldda,
dAT(j, j+1), ldda);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), ldda,
dAT(j+1, j ), ldda,
c_one, dAT(j+1, j+1), ldda );
}
else {
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(j, j ), ldda,
dAT(j, j+1), ldda);
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), ldda,
dAT(j+1, j ), ldda,
c_one, dAT(j+1, j+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_stranspose( nb0, rows, dAT(s,s), ldda, dA, cols );
magma_sgetmatrix( rows, nb0, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_sgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo);
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
for( i=s*nb; i < s*nb + nb0; ++i ) {
ipiv[i] += s*nb;
}
magmablas_slaswp( n, dAT, ldda, s*nb + 1, s*nb + nb0, ipiv, 1 );
// upload j-th panel
magma_ssetmatrix( rows, nb0, work, lda, dA, cols );
magmablas_stranspose( rows, nb0, dA, cols, dAT(s,s), ldda );
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), ldda,
dAT(s,s)+nb0, ldda);
}
// undo transpose
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_stranspose_inplace( ldda, dAT, ldda );
magma_sgetmatrix( m, n, da, ldda, A, lda );
}
else {
magmablas_stranspose( n, m, dAT, ldda, da, maxm );
magma_sgetmatrix( m, n, da, maxm, A, lda );
magma_free( dAT );
}
magma_free( dA );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_sgetrf */
/**
Purpose
-------
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf(magma_int_t m, magma_int_t n, magmaFloatComplex *A, magma_int_t lda,
magma_int_t *ipiv, magma_int_t *info)
{
#define dAT(i,j) (dAT + (i)*nb*ldda + (j)*nb)
magmaFloatComplex *dAT, *dA, *da, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
nb = magma_get_cgetrf_nb(m);
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_cgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, maxdim;
magma_int_t i, rows, cols, s = min(m, n)/nb;
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
/* set number of GPUs */
magma_int_t num_gpus = magma_num_gpus();
if ( num_gpus > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
/* explicitly checking the memory requirement */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
int h = 1+(2+num_gpus), num_gpus2 = num_gpus;
int NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
if ( num_gpus > ceil((float)NB/nb) ) {
num_gpus2 = (int)ceil((float)NB/nb);
h = 1+(2+num_gpus2);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
}
if ( num_gpus2*NB < n ) {
/* require too much memory, so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
ldda = maxn;
work = A;
if (maxdim*maxdim < 2*maxm*maxn) {
// if close to square, allocate square matrix and transpose in-place
if (MAGMA_SUCCESS != magma_cmalloc( &dA, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
ldda = maxdim;
magma_csetmatrix( m, n, A, lda, da, ldda );
dAT = da;
magmablas_ctranspose_inplace( ldda, dAT, ldda );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
magma_csetmatrix( m, n, A, lda, da, maxm );
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dA );
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
magmablas_ctranspose( m, n, da, maxm, dAT, ldda );
}
lapackf77_cgetrf( &m, &nb, work, &lda, ipiv, &iinfo);
/* Define user stream if current stream is NULL */
cudaStream_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else
stream[1] = current_stream;
for( i = 0; i < s; i++ ) {
// download i-th panel
cols = maxm - i*nb;
if (i > 0) {
// download i-th panel
magmablas_ctranspose( nb, cols, dAT(i,i), ldda, dA, cols );
// make sure that gpu queue is empty
magma_device_sync();
magma_cgetmatrix_async( m-i*nb, nb, dA, cols, work, lda,
stream[0]);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), ldda,
dAT(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), ldda,
dAT(i, i-1), ldda,
c_one, dAT(i, i+1), ldda );
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lda, ipiv+i*nb, &iinfo);
}
if (*info == 0 && iinfo > 0)
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lda, dA, cols,
stream[0]);
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
magmablas_ctranspose( cols, nb, dA, cols, dAT(i,i), ldda );
// do the small non-parallel computations
if (s > (i+1)) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), ldda,
dAT(i+1, i ), ldda,
c_one, dAT(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), ldda,
dAT(i+1, i ), ldda,
c_one, dAT(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose( nb0, rows, dAT(s,s), ldda, dA, cols );
magma_cgetmatrix( rows, nb0, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo);
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb0, s*nb );
magma_csetmatrix( rows, nb0, work, lda, dA, cols );
magmablas_ctranspose( rows, nb0, dA, cols, dAT(s,s), ldda );
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s, s), ldda,
dAT(s, s)+nb0, ldda);
}
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_ctranspose_inplace( ldda, dAT, ldda );
magma_cgetmatrix( m, n, da, ldda, A, lda );
} else {
magmablas_ctranspose( n, m, dAT, ldda, da, maxm );
magma_cgetmatrix( m, n, da, maxm, A, lda );
magma_free( dAT );
}
magma_free( dA );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* magma_cgetrf */
/**
Purpose
-------
CLAHRU is an auxiliary MAGMA routine that is used in CGEHRD to update
the trailing sub-matrices after the reductions of the corresponding
panels.
See further details below.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ihi INTEGER
Last row to update. Same as IHI in cgehrd.
@param[in]
k INTEGER
Number of rows of the matrix Am (see details below)
@param[in]
nb INTEGER
Block size
@param[out]
A COMPLEX array, dimension (LDA,N-K)
On entry, the N-by-(N-K) general matrix to be updated. The
computation is done on the GPU. After Am is updated on the GPU
only Am(1:NB) is transferred to the CPU - to update the
corresponding Am matrix. See Further Details below.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
This implementation follows the algorithm and notations described in:
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
The difference is that here Am is computed on the GPU.
M is renamed Am, G is renamed Ag.
@ingroup magma_cgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_clahru_m(
magma_int_t n, magma_int_t ihi, magma_int_t k, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
struct cgehrd_data* data )
{
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dW( d, i, j ) (data->W [d] + (i) + (j)*ldda)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t d;
magma_int_t dk, dkhi, dknb, dn;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (ihi < 0 || ihi > n) {
info = -2;
} else if (k < 0 || k > n) {
info = -3;
} else if (nb < 1 || nb > n) {
info = -4;
} else if (lda < max(1,n)) {
info = -6;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k, ihi, &dk, &dkhi );
magma_indices_1D_bcyclic( nb, ngpu, d, k+nb, n, &dknb, &dn );
// -----
// on right, A := A Q = A - A V T V'
// Update Am = Am - Am V T Vd' = Am - Ym Wd', with Wd = Vd T'
// Wd = Vd T' = V(k:ihi-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
// Vd and Wd are the portions corresponding to the block cyclic dkstribution
magma_cgemm( MagmaNoTrans, MagmaConjTrans, dkhi-dk, nb, nb,
c_one, dVd(d, dk, 0), ldvd,
dTi(d), nb,
c_zero, dW (d, dk, 0), ldda );
// Am = Am - Ym Wd' = A(0:k-1, k:ihi-1) - Ym(0:k-1, 0:nb-1) * W(k:ihi-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, k, dkhi-dk, nb,
c_neg_one, dY(d, 0, 0), ldda,
dW(d, dk, 0), ldda,
c_one, dA(d, 0, dk), ldda );
// -----
// on right, A := A Q = A - A V T V'
// Update Ag = Ag - Ag V T V' = Ag - Yg Wd'
// Ag = Ag - Yg Wd' = A(k:ihi-1, nb:ihi-k-1) - Y(k:ihi-1, 0:nb-1) * W(k+nb:ihi-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, dkhi-dknb, nb,
c_neg_one, dY(d, k, 0), ldda,
dW(d, dknb, 0), ldda,
c_one, dA(d, k, dknb), ldda );
// -----
// on left, A := Q' A = A - V T' V' A
// Ag2 = Ag2 - V T' V' Ag2 = W Yg, with W = V T' and Yg = V' Ag2
// Note that Ag is A(k:ihi, nb+1:ihi-k)
// while Ag2 is A(k:ihi, nb+1: n -k)
// here V and W are the whole matrices, not just block cyclic portion
// W = V T' = V(k:ihi-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
// TODO would it be cheaper to compute the whole matrix and
// copy the block cyclic portions to another workspace?
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, nb, nb,
c_one, dV (d, k, 0), ldv,
dTi(d), nb,
c_zero, dW (d, k, 0), ldda );
// Z = V(k:ihi-1, 0:nb-1)' * A(k:ihi-1, nb:n-k-1); Z is stored over Y
magma_cgemm( MagmaConjTrans, MagmaNoTrans, nb, dn-dknb, ihi-k,
c_one, dV(d, k, 0), ldv,
dA(d, k, dknb), ldda,
c_zero, dY(d, 0, 0), nb );
// Ag2 = Ag2 - W Z = A(k:ihi-1, k+nb:n-1) - W(k+nb:n-1, 0:nb-1) * Z(0:nb-1, k+nb:n-1)
magma_cgemm( MagmaNoTrans, MagmaNoTrans, ihi-k, dn-dknb, nb,
c_neg_one, dW(d, k, 0), ldda,
dY(d, 0, 0), nb,
c_one, dA(d, k, dknb), ldda );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return info;
}
/**
Purpose
-------
SSYTRD reduces a real symmetric matrix A to real symmetric
tridiagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
Arguments
---------
@param[in]
num_gpus INTEGER
The number of GPUs. num_gpus > 0.
@param[in]
num_streams INTEGER
The number of GPU streams used for update. 10 >= num_streams > 0.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = MagmaUpper, the diagonal and first superdiagonal
of A are overwritten by the corresponding elements of the
tridiagonal matrix T, and the elements above the first
superdiagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the diagonal and first subdiagonal of A are over-
written by the corresponding elements of the tridiagonal
matrix T, and the elements below the first subdiagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
d REAL array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e REAL array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
tau REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_ssyev_comp
********************************************************************/
extern "C" magma_int_t
magma_ssytrd_mgpu(
magma_int_t num_gpus, magma_int_t num_streams, magma_uplo_t uplo, magma_int_t n,
float *A, magma_int_t lda,
float *d, float *e, float *tau,
float *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i, j) (A + (j)*lda + (i))
#define dA(id, i, j) (dA[(id)] + (j)*ldda + (i))
#define dW(id, i, j) (dwork[(id)] + (j)*ldda + (i))
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t ln, ldda;
magma_int_t nb = magma_get_ssytrd_nb(n), ib;
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
float d_one = MAGMA_D_ONE;
//float mv_time = 0.0;
#ifdef PROFILE_SY2RK
float up_time = 0.0;
#endif
magma_int_t kk, nx;
magma_int_t i = 0, ii, iii, j, did, i_n;
magma_int_t iinfo;
magma_int_t ldwork, lddwork, lwkopt, ldwork2;
magma_int_t lquery;
magma_queue_t stream[MagmaMaxGPUs][10];
float *dx[MagmaMaxGPUs], *dy[MagmaMaxGPUs], *hwork;
float *dwork2[MagmaMaxGPUs];
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < nb*n && ! lquery) {
*info = -9;
} else if ( num_streams > 2 ) {
*info = 2; // TODO fix
}
/* Determine the block size. */
ldwork = lddwork = n;
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_S_MAKE( lwkopt, 0 );
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
float *dA[MagmaMaxGPUs];
float *dwork[MagmaMaxGPUs];
float times[11];
for( did=0; did < 11; did++ )
times[did] = 0;
//#define PROFILE_SY2RK
#ifdef PROFILE_SY2RK
magma_event_t start, stop;
float etime;
magma_setdevice(0);
magma_event_create( &start );
magma_event_create( &stop );
#endif
ldda = lda;
ln = ((nb*(1+n/(nb*num_gpus))+31)/32)*32;
ldwork2 = (1+ n / nb + (n % nb != 0)) * ldda;
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
// TODO fix memory leak
if ( MAGMA_SUCCESS != magma_smalloc(&dA[did], ln*ldda+3*lddwork*nb) ||
MAGMA_SUCCESS != magma_smalloc(&dx[did], num_streams*n) ||
MAGMA_SUCCESS != magma_smalloc(&dy[did], num_streams*n) ||
MAGMA_SUCCESS != magma_smalloc(&dwork2[did], ldwork2 ) ) {
for( i=0; i < did; i++ ) {
magma_setdevice(i);
magma_free(dA[i]);
magma_free(dx[i]);
magma_free(dy[i]);
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork[did] = dA[did] + ln*ldda;
for( kk=0; kk < num_streams; kk++ )
magma_queue_create(&stream[did][kk]);
}
magma_setdevice(0);
// TODO fix memory leak dwork2
if ( MAGMA_SUCCESS != magma_smalloc_pinned( &hwork, num_streams*num_gpus*n ) ) {
for( i=0; i < num_gpus; i++ ) {
magma_setdevice(i);
magma_free(dA[i]);
magma_free(dx[i]);
magma_free(dy[i]);
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if (n < 2048)
nx = n;
else
nx = 512;
if (upper) {
/* Copy the matrix to the GPU */
if (1 <= n-nx) {
magma_shtodhe(num_gpus, uplo, n, nb, A, lda, dA, ldda, stream, &iinfo );
}
/* Reduce the upper triangle of A.
Columns 1:kk are handled by the unblocked method. */
for (i = nb*((n-1)/nb); i >= nx; i -= nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
/* wait for the next panel */
if (i != nb*((n-1)/nb)) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
magma_slatrd_mgpu(num_gpus, uplo, n, i+ib, ib, nb,
A(0, 0), lda, e, tau,
work, ldwork,
dA, ldda, 0,
dwork, i+ib,
dwork2, ldwork2,
1, dx, dy, hwork,
stream, times);
magma_ssyr2k_mgpu(num_gpus, MagmaUpper, MagmaNoTrans, nb, i, ib,
c_neg_one, dwork, i+ib, 0,
d_one, dA, ldda, 0,
num_streams, stream);
/* get the next panel */
if (i-nb >= nx ) {
ib = min(nb, n-(i-nb));
ii = nb*((i-nb)/(nb*num_gpus));
did = ((i-nb)/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( (i-nb)+ib, ib,
dA(did, 0, ii), ldda,
A(0, i-nb), lda,
stream[did][0] );
}
/* Copy superdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+ib; ++j) {
if ( j > 0 ) {
*A(j-1,j) = MAGMA_S_MAKE( e[j - 1], 0 );
}
d[j] = MAGMA_S_REAL( *A(j, j) );
}
} /* end of for i=... */
if ( nx > 0 ) {
if (1 <= n-nx) { /* else A is already on CPU */
for (i=0; i < nx; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( nx, ib,
dA(did, 0, ii), ldda,
A(0, i), lda,
stream[did][0] );
}
}
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
/* Use unblocked code to reduce the last or only block */
lapackf77_ssytd2(uplo_, &nx, A(0, 0), &lda, d, e, tau, &iinfo);
}
}
else {
trace_init( 1, num_gpus, num_streams, (CUstream_st**)stream );
/* Copy the matrix to the GPU */
if (1 <= n-nx) {
magma_shtodhe(num_gpus, uplo, n, nb, A, lda, dA, ldda, stream, &iinfo );
}
/* Reduce the lower triangle of A */
for (i = 0; i < n-nx; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
/* Reduce columns i:i+ib-1 to tridiagonal form and form the
matrix W which is needed to update the unreduced part of
the matrix */
/* Get the current panel (no need for the 1st iteration) */
if (i != 0) {
magma_setdevice(did);
trace_gpu_start( did, 0, "comm", "get" );
magma_sgetmatrix_async( n-i, ib,
dA(did, i, ii), ldda,
A(i,i), lda,
stream[did][0] );
trace_gpu_end( did, 0 );
magma_queue_sync(stream[did][0]);
magma_setdevice(0);
}
magma_slatrd_mgpu(num_gpus, uplo, n, n-i, ib, nb,
A(i, i), lda, &e[i],
&tau[i], work, ldwork,
dA, ldda, i,
dwork, (n-i),
dwork2, ldwork2,
1, dx, dy, hwork,
stream, times );
#ifdef PROFILE_SY2RK
magma_setdevice(0);
if ( i > 0 ) {
cudaEventElapsedTime(&etime, start, stop);
up_time += (etime/1000.0);
}
magma_event_record(start, 0);
#endif
magma_ssyr2k_mgpu(num_gpus, MagmaLower, MagmaNoTrans, nb, n-i-ib, ib,
c_neg_one, dwork, n-i, ib,
d_one, dA, ldda, i+ib, num_streams, stream);
#ifdef PROFILE_SY2RK
magma_setdevice(0);
magma_event_record(stop, 0);
#endif
/* Copy subdiagonal elements back into A, and diagonal
elements into D */
for (j = i; j < i+ib; ++j) {
if ( j+1 < n ) {
*A(j+1,j) = MAGMA_S_MAKE( e[j], 0 );
}
d[j] = MAGMA_S_REAL( *A(j, j) );
}
} /* for i=... */
/* Use unblocked code to reduce the last or only block */
if ( i < n ) {
iii = i;
i_n = n-i;
if ( i > 0 ) {
for (; i < n; i += nb) {
ib = min(nb, n-i);
ii = nb*(i/(nb*num_gpus));
did = (i/nb)%num_gpus;
magma_setdevice(did);
magma_sgetmatrix_async( i_n, ib,
dA(did, iii, ii), ldda,
A(iii, i), lda,
stream[did][0] );
}
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
magma_queue_sync(stream[did][0]);
}
}
lapackf77_ssytrd(uplo_, &i_n, A(iii, iii), &lda, &d[iii], &e[iii],
&tau[iii], work, &lwork, &iinfo);
}
}
#ifdef PROFILE_SY2RK
magma_setdevice(0);
if ( n > nx ) {
cudaEventElapsedTime(&etime, start, stop);
up_time += (etime/1000.0);
}
magma_event_destroy( start );
magma_event_destroy( stop );
#endif
trace_finalize( "ssytrd.svg", "trace.css" );
for( did=0; did < num_gpus; did++ ) {
magma_setdevice(did);
for( kk=0; kk < num_streams; kk++ )
magma_queue_sync(stream[did][kk]);
for( kk=0; kk < num_streams; kk++ )
magma_queue_destroy(stream[did][kk]);
magma_free(dA[did]);
magma_free(dx[did]);
magma_free(dy[did]);
magma_free(dwork2[did]);
}
magma_free_pinned(hwork);
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
work[0] = MAGMA_S_MAKE( lwkopt, 0 );
#ifdef PROFILE_SY2RK
printf( " n=%d nb=%d\n", n, nb );
printf( " Time in SLARFG: %.2e seconds\n", times[0] );
//printf( " Time in SSYMV : %.2e seconds\n", mv_time );
printf( " Time in SSYR2K: %.2e seconds\n", up_time );
#endif
return *info;
} /* magma_ssytrd */
extern "C" magma_int_t
magma_zbulge_applyQ_v2_m(
magma_int_t ngpu,
magma_side_t side,
magma_int_t NE, magma_int_t N,
magma_int_t NB, magma_int_t Vblksiz,
magmaDoubleComplex *E, magma_int_t lde,
magmaDoubleComplex *V, magma_int_t ldv,
magmaDoubleComplex *T, magma_int_t ldt,
magma_int_t *info)
{
//%===========================
//% local variables
//%===========================
magma_int_t Vm, Vn, mt, nt;
magma_int_t myrow, mycol, blkj, blki;
magma_int_t blkid,vpos,tpos;
magma_int_t firstrow, nbcolinvolvd;
magma_int_t versionL = 113;
magma_int_t versionR = 92;
magma_int_t Vchunksiz = 10;
*info=0;
/* Quick return */
if ( NE == 0 ) {
return *info;
}
if ( N == 0 ) {
return *info;
}
if ( NB == 0 ) {
return *info;
}
/* ==========================================
* some infos for developer
* Initialisation and checking nb of cores
* ==========================================*/
/* we have 2 algo for left (113 114) and 2 algo for right (91 92)
* which correspond to versionL versionR.
* They are very similar (detail explained in tech report and matlab code)
* however version 114 and 92 improve locality.
* while version 113 is used in case WNATZ=1 (construct Q2) which allow
* the construction to be done in an optimized way taking into
* consideration that the matrix is Identity so making less flops.
*
*/
// Initialize streaming and events
magma_device_sync();
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_int_t nbevents =2, nstream=2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t myevent[MagmaMaxGPUs][20];
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[dev][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&myevent[dev][i],cudaEventDisableTiming);
}
}
// Azzam 21/11/2012
// NOTE THAT dwork was of size 2*NE*Vblksiz+...
// but I am thinking why not modifing it to NE*Vblksiz+...
// BUT NO because the 2* is used because of making 2 streams working and so
// they might be using dwork in parallel
magmaDoubleComplex *dE[MagmaMaxGPUs];
magmaDoubleComplex *dwork[MagmaMaxGPUs], *dwork0[MagmaMaxGPUs], *dwork1[MagmaMaxGPUs];
//magmaDoubleComplex *dwvt[MagmaMaxGPUs];
magmaDoubleComplex *dwvt0[MagmaMaxGPUs], *dwvt1[MagmaMaxGPUs];
magmaDoubleComplex *dT0[MagmaMaxGPUs], *dV0[MagmaMaxGPUs], *dT1[MagmaMaxGPUs], *dV1[MagmaMaxGPUs];
magma_int_t dev;
magma_int_t ldde = N;
magma_int_t lddv = ldv;
magma_int_t lddt = ldt;
magma_int_t ne_loc = magma_ceildiv(NE, ngpu);
if (ne_loc < 256)
ne_loc=256;
magma_int_t dwVTsiz = lddv*Vblksiz; // lddv*lddv + lddv*NE; // lddv*Vblksiz;
magma_int_t dworksiz = ne_loc*Vblksiz; // lddv*Vblksiz; // NE*Vblksiz;
ngpu = min(ngpu, magma_ceildiv(NE,ne_loc)); // Don't use GPU that will not have data.
// copy dE to GPUs
for (dev=0; dev < ngpu; ++dev) {
magma_setdevice( dev );
if (MAGMA_SUCCESS != magma_zmalloc( &dE[dev], ldde * ne_loc)) {
printf ("!!!! magma_zbulge_applyQ magma_alloc failed for: dE\n" );
exit(-1);
}
if (MAGMA_SUCCESS != magma_zmalloc( &dwork[dev], 2*dworksiz + 2*dwVTsiz + 2*Vchunksiz* (Vblksiz* (lddv+lddt)) )) {
printf ("!!!! magma_zbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
dwork0[dev] = dwork[dev]; // size = dworksiz;
dwork1[dev] = dwork0[dev] + dworksiz; // size = dworksiz;
dwvt0[dev] = dwork[dev] + 2*dworksiz; // size = dwVTsiz;
dwvt1[dev] = dwvt0[dev] + dwVTsiz; // size = dwVTsiz;
dV0[dev] = dwork[dev] + 2*dworksiz + 2*dwVTsiz;
dT0[dev] = dV0[dev] + Vchunksiz*Vblksiz*lddv;
dV1[dev] = dT0[dev] + Vchunksiz*Vblksiz*lddt;
dT1[dev] = dV1[dev] + Vchunksiz*Vblksiz*lddv;
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_zsetmatrix_async( N, ie_loc, E+lde*ne_loc*dev, lde, dE(dev, 0, 0), ldde, streams[dev][1] );
}
// make overlapped copy
magma_int_t ncpy = 0;
magma_int_t copyed=0, copyst=0;
magma_int_t blkcnt,nothing, mysiz, flip, vld,tld, locpos;
findVTsiz(N, NB, Vblksiz, &blkcnt, ¬hing);
flip = 0;
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
#ifdef ENABLE_DEBUG
printf(" APPLY Q_v22_m GPU with NGPU %d N %d, NE %d, NB %d, Vblksiz %d, versionL %d versionR %d SIDE %c \n",
ngpu, N, NE, NB, Vblksiz, versionL, versionR, side);
#endif
/*
* MagmamaLeft
*/
if (side == MagmaLeft) {
/*
* Version 113:
* loop over the block_col (nt) and for each find the
* number of tiles (mt) in this block_col. then loop over mt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
if ( versionL == 113 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=nt-1; blkj >= 0; blkj--) {
/* the index of the first row on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=mt; blki > 0; blki--) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d vpos %d \n",blkj,blki,Vm, Vn,mycol,vpos);
magma_bulge_findpos113(N, NB, Vblksiz, mycol, myrow, &blkid);
// COPY Vchunksiz Vs and Vchunksiz Ts to GPU and store it in dV0/dV1 and dT0/dT1
if (ncpy == 0) {
// flip = 1 for this.
copyst = 0; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
if (mysiz > 0) {
ncpy = 1;
flip = 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_zsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, streams[dev][1]);
magma_zsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, streams[dev][1]);
}
//printf("doing the first copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
}
}
if (blkid == copyst) {
flip = ncpy % 2;
copyst = copyed; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
//printf(" get to copy blkid %d blkid+(2*Vchunksiz) %d copyst %d copyed %d\n",blkid,blkid+(Vchunksiz),copyst,copyed);
if (mysiz > 0) {
ncpy = ncpy + 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
if (flip == 0) { // now I am working on dV0 so copy the next and put it on dV1
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_zsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, streams[dev][1]);
magma_zsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, streams[dev][1]);
}
} else { // now I am working on dV1 so copy the next and put it on dV0
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV0 dT0\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_zsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV0[dev], vld, streams[dev][0]);
magma_zsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT0[dev], tld, streams[dev][0]);
}
}
}
}
if ((Vm > 0) && (Vn > 0)) {
locpos = blkid%Vchunksiz;
magma_int_t lcvpos = locpos*Vblksiz*lddv;
magma_int_t lctpos = locpos*Vblksiz*lddt;
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d locvpos %5d loctpos %5d blkid %2d using data in dV%1d dT%1d \n",blkj,blki,Vm, Vn,mycol,lcvpos,lctpos, blkid,flip,flip);
if (flip == 0) {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
magma_queue_wait_event( streams[dev][0], myevent[dev][1] );
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib);
magma_zlarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib, dwvt0[dev], Vm);
}
magma_event_record( myevent[dev][0], streams[dev][0] );
}
} else {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][1]);
magma_queue_wait_event( streams[dev][1], myevent[dev][0] );
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib);
magma_zlarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib, dwvt1[dev], Vm);
}
magma_event_record( myevent[dev][1], streams[dev][1] );
}
}
} // end for (Vm &Vn) > 0
} // end for blki
} // end for blkj
} // end if version=113
/*
* Version 114:
* loop over the block_row (mt) and for each find diagonally the
* number of tiles (nt) in this block_row. then loop over nt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
else {
printf("versionL 114 not implemented in zbulge_applyQ_v2_m\n");
exit(-1);
mt = magma_ceildiv((N-1),NB);
for (blki = mt; blki > 0; blki--) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = nt-1; blkj >= 0; blkj--) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_zsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_zsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
//printf("voici blki %d rownbm %d mycol %d coled %d blkid %d vpos %d tpos %d\n", blki, rownbm, mycol, coled, blkid, vpos, tpos);
for (magma_int_t i=0; i < NE; i += sz_bl) {
ib = min(sz_bl, NE-i);
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0, lddv, dT0, lddt, dE(myrow,i), ldde, dwork, NE);
}
*/
} // end for (Vm &Vn) > 0
} // end for blkj
} // end for blki
} // end version 114
} // end LEFT
/*
* MagmaRight
*/
else {
printf("Side 'R' not implemented in zbulge_applyQ_v2_m\n");
exit(-1);
/*
* Version 91:
*/
if ( versionR == 91 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=0; blkj < nt; blkj++) {
/* the index of the first myrow on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=1; blki <= mt; blki++) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( (blkj == nt-1) && (blki == mt) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
mycol = blkj*Vblksiz;
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_zsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_zsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
*/
} // end for (Vm &Vn) > 0
} // end for blki
} // end fo blkj
} // end of version 91
/*
* Version 92:
*/
else {
mt = magma_ceildiv((N-1),NB);
for (blki = 1; blki <= mt; blki++) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = 0; blkj < nt; blkj++) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_zsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_zsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
*/
} // end for (Vm &Vn) > 0
} //end for blkj
} // end for blki
} //end of version 92
} // end RIGHT
// copy back the dE form each GPU
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
magma_queue_wait_event( streams[dev][0], myevent[dev][1] );
magma_queue_wait_event( streams[dev][0], myevent[dev][0] );
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_zgetmatrix_async( N, ie_loc, dE(dev, 0, 0), ldde, E+lde*ne_loc*dev, lde, streams[dev][0] );
magma_event_record( myevent[dev][0], streams[dev][0] );
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
magma_queue_wait_event( streams[dev][0], myevent[dev][0] );
magma_device_sync(); // no need for synchronize
magma_free(dwork[dev]);
magma_free(dE[dev]);
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( myevent[dev][i] );
}
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[dev][i] );
}
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
}
extern "C" void
magma_ssyr2k_mgpu(
magma_int_t num_gpus, magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
float alpha,
float **db, magma_int_t lddb, magma_int_t offset_b,
float beta,
float **dc, magma_int_t lddc, magma_int_t offset,
magma_int_t num_streams, magma_queue_t stream[][10])
{
#define dB(id, i, j) (db[(id)]+(j)*lddb + (i)+offset_b)
#define dB1(id, i, j) (db[(id)]+(j)*lddb + (i)+offset_b)+k*lddb
#define dC(id, i, j) (dc[(id)]+(j)*lddc + (i))
magma_int_t i, id, ib, ii, kk, n1;
float c_one = MAGMA_S_ONE;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
/* ssyr2k on diagonal block */
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_ssyr2k(uplo, trans, ib, k,
alpha, dB1(id, i, 0 ), lddb,
dB(id, i, 0 ), lddb,
beta, dC(id, i+offset, ii), lddc);
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
magma_sgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB1(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc);
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
n1 = n-i-ib;
// sgemm on off-diagonal blocks
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_sgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB1(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, i+offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
magma_sgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB( id, 0, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc);
}
} else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
n1 = n-i-ib;
/* sgemm on off-diagonal blocks */
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_sgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB(id, i+ib, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, i+offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
for( id=0; id < num_gpus; id++ ) {
magma_setdevice(id);
for( kk=0; kk < num_streams; kk++ ) {
magma_queue_sync(stream[id][kk]);
}
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
