extern "C" magma_int_t
magmablas_dsymv_sync( magma_int_t num_gpus, magma_int_t k,
magma_int_t n, double *work, double *W,
magma_queue_t stream[][10] )
{
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
magma_int_t id, kk;
/* reduce on CPU */
magma_setdevice(0);
magma_queue_sync(stream[0][0]);
for( kk=1; kk < k; kk++ ) {
magma_queue_sync(stream[0][kk]);
blasf77_daxpy( &n, &c_one, &work[kk*n], &ione, W, &ione );
}
for( id=1; id < num_gpus; id++ ) {
magma_setdevice(id);
for( kk=0; kk < k; kk++ ) {
magma_queue_sync(stream[id][kk]);
blasf77_daxpy( &n, &c_one, &work[id*k*n + kk*n], &ione, W, &ione );
}
}
return 0;
}
magma_int_t
magmablas_zhemv_sync( magma_int_t num_gpus, magma_int_t k,
magma_int_t n, magmaDoubleComplex *work, magmaDoubleComplex *w,
magma_queue_t stream[][10] )
{
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magma_int_t id, kk;
/* reduce on CPU */
magma_setdevice(0);
magma_queue_sync(stream[0][0]);
for( kk=1; kk<k; kk++ ) {
magma_queue_sync(stream[0][kk]);
blasf77_zaxpy( &n, &c_one, &work[kk*n], &ione, w, &ione );
}
for( id=1; id<num_gpus; id++ ) {
magma_setdevice(id);
for( kk=0; kk<k; kk++ ) {
magma_queue_sync(stream[id][kk]);
blasf77_zaxpy( &n, &c_one, &work[id*k*n + kk*n], &ione, w, &ione );
}
}
return 0;
}
extern "C" magma_int_t
magma_zdtohpo(magma_int_t num_gpus, char *uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb, magma_int_t NB,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *dwork[], magma_int_t ldda,
magma_queue_t stream[][3], magma_int_t *info)
{
magma_int_t k;
if( lapackf77_lsame(uplo, "U") ) {
magma_int_t j, jj, jb, mj;
/* go through each column */
for (j=off_j+NB; j<n; j+=nb) {
jj = (j-off_j)/(nb*num_gpus);
k = ((j-off_j)/nb)%num_gpus;
jb = min(nb, (n-j));
if(j+jb < off_j+m)
mj = (j-off_i)+jb;
else
mj = m;
magma_setdevice(k);
magma_zgetmatrix_async( mj, jb,
dA(k, 0, jj*nb), ldda,
A(off_i, j), lda,
stream[k][0] );
magma_queue_sync( stream[k][0] );
}
} else {
magma_int_t i, ii, ib, ni;
/* go through each row */
for(i=off_i+NB; i<m; i+=nb) {
ii = (i-off_i)/(nb*num_gpus);
k = ((i-off_i)/nb)%num_gpus;
ib = min(nb, (m-i));
if(i+ib < off_i+n)
ni = (i-off_i)+ib;
else
ni = n;
magma_setdevice(k);
magma_zgetmatrix_async( ib, ni,
dA(k, ii*nb, 0), ldda,
A(i, off_j), lda,
stream[k][0] );
magma_queue_sync( stream[k][0] );
}
}
/*for( k=0; k<num_gpus; k++ ) {
magma_setdevice(k);
magma_queue_sync( stream[k][0] );
}*/
magma_setdevice(0);
return *info;
}
void magma_task_dev_dgetmatrix_async_transpose(Schedule* sched_obj )
{
magma_int_t deviceID;
magma_int_t m;
magma_int_t nb;
double *dA_src;
magma_int_t dA_LD;
double *A_dst;
magma_int_t LDA;
magma_queue_t stream1;
double *dwork;
magma_int_t dwork_LD;
void *dep_ptr;
#if (dbglevel >=1)
ca_trace_start();
#endif
// printf("Matrix_get_transpose\n");
//schedule_unpack_args_8(sched_obj, m, nb, dA_src, dA_LD, A_dst, LDA, dwork, dwork_LD);
schedule_unpack_args_11(sched_obj, deviceID, m, nb, dA_src, dA_LD, A_dst, LDA, stream1, dwork, dwork_LD, dep_ptr);
magma_setdevice(deviceID);
// printf("Matrix_get_transpose m:%d, nb:%d, dep_ptr:%p\n",m,nb,dep_ptr);
//magma_dgetmatrix_transpose( m, nb, dA_src, dA_LD, A_dst, LDA);
// task_getpanel(gpu_nrows, ncols, dAT(K,K+A_N), dAT_LD, A(K,K+A_N), A_LD);
#if (dbglevel==10)
ca_dbg_printMat_transpose_gpu(nb, m, dA_src, dA_LD, "dA before getMatrix");
#endif
// pthread_mutex_lock(&mutex_dAP);
/*1.transpose dA to dwork*/
magmablas_dtranspose2s(dwork, dwork_LD, dA_src, dA_LD, nb, m, compute_stream[deviceID]);
// printf("Matrix_get_transpose m:%d, nb:%d 1:done\n",m,nb);
#if (dbglevel==10)
// ca_dbg_printMat_gpu(m, nb, dwork, dwork_LD, "dwork after dA transpose");
#endif
magma_queue_sync(compute_stream[deviceID]);
/*2. copy dwork to A: send the panel to GPU*/
magma_dgetmatrix_async(m, nb, dwork, dwork_LD, A_dst, LDA, stream1);
// printf("Matrix_get_transpose m:%d, nb:%d 2:done\n",m,nb);
// pthread_mutex_unlock(&mutex_dAP);
magma_queue_sync(stream1); //make sure dwork is set before the transpose
#if (dbglevel==10)
ca_dbg_printMat(m, nb, A_dst, LDA, "A after getMatrix");
#endif
#if (dbglevel >=1)
ca_trace_end_gpu('G');
ca_trace_end_cpu('C');
#endif
// printf("End Matrix_get_transpose m:%d, nb:%d\n",m,nb);
}
// ----------------------------------------------------------------------
// TODO info is unused
extern "C" magma_int_t
magma_zhtodhe(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaDoubleComplex *A, magma_int_t lda,
magmaDoubleComplex_ptr dA[], magma_int_t ldda,
magma_queue_t queues[][10],
magma_int_t *info)
{
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_int_t k;
if (uplo == MagmaLower) {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*ngpu);
k = (j/nb)%ngpu;
jb = min(nb, (n-j));
mj = n-j;
magma_setdevice( k );
magma_zsetmatrix_async( mj, jb,
A(j,j), lda,
dA(k, j, jj*nb), ldda,
queues[k][0] );
}
}
else {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*ngpu);
k = (j/nb)%ngpu;
jb = min(nb, (n-j));
mj = j+jb;
magma_setdevice( k );
magma_zsetmatrix_async( mj, jb,
A(0, j), lda,
dA(k, 0, jj*nb), ldda,
queues[k][0] );
}
}
for( k=0; k < ngpu; k++ ) {
magma_setdevice( k );
magma_queue_sync( queues[k][0] );
}
magma_setdevice( orig_dev );
return *info;
}
extern "C" magma_int_t
magma_shtodhe(magma_int_t num_gpus, magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
float *A, magma_int_t lda,
float **dA, magma_int_t ldda,
magma_queue_t stream[][10], magma_int_t *info)
{
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_int_t k;
if (uplo == MagmaLower) {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*num_gpus);
k = (j/nb)%num_gpus;
jb = min(nb, (n-j));
mj = n-j;
magma_setdevice(k);
magma_ssetmatrix_async( mj, jb,
A(j,j), lda,
dA(k, j, jj*nb), ldda,
stream[k][0] );
}
}
else {
/* go through each block-column */
magma_int_t j, jj, jb, mj;
for (j=0; j < n; j += nb) {
jj = j/(nb*num_gpus);
k = (j/nb)%num_gpus;
jb = min(nb, (n-j));
mj = j+jb;
magma_setdevice(k);
magma_ssetmatrix_async( mj, jb,
A(0, j), lda,
dA(k, 0, jj*nb), ldda,
stream[k][0] );
}
}
for( k=0; k < num_gpus; k++ ) {
magma_setdevice(k);
magma_queue_sync(stream[k][0]);
}
magma_setdevice( orig_dev );
return *info;
}
extern "C" magma_int_t
magma_spotrf_batched(
magma_uplo_t uplo, magma_int_t n,
float **dA_array, magma_int_t ldda,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower) {
arginfo = -1;
} else if (n < 0) {
arginfo = -2;
} else if (ldda < max(1,n)) {
arginfo = -4;
}
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
// Quick return if possible
if (n == 0) {
return arginfo;
}
magma_int_t crossover = magma_get_spotrf_batched_crossover();
if (n > crossover )
{
// The memory allocation/deallocation inside this routine takes about 3-4ms
arginfo = magma_spotrf_lg_batched(uplo, n, dA_array, ldda, info_array, batchCount, queue);
}
else
{
#if defined(VERSION20)
arginfo = magma_spotrf_lpout_batched(uplo, n, dA_array, ldda, 0, info_array, batchCount, queue);
#elif defined(VERSION33)
arginfo = magma_spotrf_v33_batched(uplo, n, dA_array, ldda, info_array, batchCount, queue);
#elif defined(VERSION31)
arginfo = magma_spotrf_lpin_batched(uplo, n, dA_array, ldda, 0, info_array, batchCount, queue);
#else
printf("ERROR NO VERSION CHOSEN\n");
#endif
}
magma_queue_sync(queue);
return arginfo;
}
magma_err_t
magma_stranspose_host( int flag, int d,
magmaFloat_ptr odata, int offseto, int ldo,
magmaFloat_ptr idata, int offseti, int ldi,
int m, int n,
magma_queue_t queue )
{
int i, j;
float *work1, *work2;
if (MAGMA_SUCCESS != magma_smalloc_cpu( &work1, m*n )) {
return MAGMA_ERR_HOST_ALLOC;
}
if (MAGMA_SUCCESS != magma_smalloc_cpu( &work2, m*n )) {
return MAGMA_ERR_HOST_ALLOC;
}
/* download to CPU */
magma_sgetmatrix_async( m, n,
idata, offseti, ldi,
work1, 0, m, queue, NULL );
magma_queue_sync( queue );
/* transpose it on CPU */
for( i=0; i<m; i++ ) {
for( j=0; j<n; j++ ) work2[j+i*n] = work1[i+j*m];
}
/* upload to GPU */
magma_ssetmatrix_async( n, m,
work2, 0, n,
odata, offseto, ldo,
queue, NULL );
magma_queue_sync( queue );
magma_free_cpu( work2 );
magma_free_cpu( work1 );
return 0;
}
void magma_task_dev_queue_sync(Schedule* sched_obj)
{
magma_int_t deviceID;
magma_queue_t stream1;
void *dep_ptr;
// double *tmpdA;
#if (dbglevel >=1)
ca_trace_start();
#endif
schedule_unpack_args_3(sched_obj, deviceID, stream1, dep_ptr);
magma_setdevice(deviceID);
magma_queue_sync(stream1);
#if (dbglevel >=1)
ca_trace_end_gpu('O');
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
-------
CGEQRF_OOC computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
This is an out-of-core (ooc) version that is similar to magma_cgeqrf but
the difference is that this version can use a GPU even if the matrix
does not fit into the GPU memory at once.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb( M, N ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgeqrf_ooc(
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
#define A(i_,j_) ( A + (i_) + (j_)*lda )
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
/* Constants */
const magmaFloatComplex c_one = MAGMA_C_ONE;
/* Local variables */
magmaFloatComplex_ptr dA, dwork;
magma_int_t i, ib, IB, j, min_mn, lddwork, ldda, rows;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
magma_int_t lwkopt = n * nb;
work[0] = magma_cmake_lwork( lwkopt );
bool lquery = (lwork == -1);
*info = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,n) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
return *info;
}
/* Check how much memory do we have */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
magma_int_t NB = magma_int_t(0.8*freeMem/m);
NB = (NB / nb) * nb;
if (NB >= n)
return magma_cgeqrf(m, n, A, lda, tau, work, lwork, info);
min_mn = min(m,n);
if (min_mn == 0) {
work[0] = c_one;
return *info;
}
lddwork = magma_roundup( NB, 32 ) + nb;
ldda = magma_roundup( m, 32 );
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (NB + nb)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magmaFloatComplex_ptr ptr = dA + ldda*NB;
dwork = dA + ldda*(NB + nb);
/* start the main loop over the blocks that fit in the GPU memory */
for (i=0; i < n; i += NB) {
IB = min( n-i, NB );
//printf("Processing %5d columns -- %5d to %5d ... \n", IB, i, i+IB);
/* 1. Copy the next part of the matrix to the GPU */
magma_csetmatrix_async( m, IB,
A(0,i), lda,
dA(0,0), ldda, queues[0] );
magma_queue_sync( queues[0] );
/* 2. Update it with the previous transformations */
for (j=0; j < min(i,min_mn); j += nb) {
ib = min( min_mn-j, nb );
/* Get a panel in ptr. */
// 1. Form the triangular factor of the block reflector
// 2. Send it to the GPU.
// 3. Put 0s in the upper triangular part of V.
// 4. Send V to the GPU in ptr.
// 5. Update the matrix.
// 6. Restore the upper part of V.
rows = m-j;
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, A(j,j), &lda, tau+j, work, &ib);
magma_csetmatrix_async( ib, ib,
work, ib,
dwork, lddwork, queues[1] );
magma_cpanel_to_q( MagmaUpper, ib, A(j,j), lda, work+ib*ib );
magma_csetmatrix_async( rows, ib,
A(j,j), lda,
ptr, rows, queues[1] );
magma_queue_sync( queues[1] );
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, IB, ib,
ptr, rows, dwork, lddwork,
dA(j, 0), ldda, dwork+ib, lddwork, queues[1] );
magma_cq_to_panel( MagmaUpper, ib, A(j,j), lda, work+ib*ib );
}
/* 3. Do a QR on the current part */
if (i < min_mn)
magma_cgeqrf2_gpu( m-i, IB, dA(i,0), ldda, tau+i, info );
/* 4. Copy the current part back to the CPU */
magma_cgetmatrix_async( m, IB,
dA(0,0), ldda,
A(0,i), lda, queues[0] );
}
magma_queue_sync( queues[0] );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_cgeqrf_ooc */
/**
Purpose
-------
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
-------
DGEQLF computes a QL factorization of a DOUBLE_PRECISION M-by-N matrix A:
A = Q * L.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A DOUBLE_PRECISION array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, if m >= n, the lower triangle of the subarray
A(m-n+1:m,1:n) contains the N-by-N lower triangular matrix L;
if m <= n, the elements on and below the (n-m)-th
superdiagonal contain the M-by-N lower trapezoidal matrix L;
the remaining elements, with the array TAU, represent the
orthogonal matrix Q as a product of elementary reflectors
(see Further Details).
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
tau DOUBLE_PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
\n
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= max(1,N,2*NB^2).
For optimum performance LWORK >= max(N*NB, 2*NB^2) where NB can be obtained
through magma_get_dgeqlf_nb(M).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(k) . . . H(2) H(1), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(m-k+i+1:m) = 0 and v(m-k+i) = 1; v(1:m-k+i-1) is stored on exit in
A(1:m-k+i-1,n-k+i), and tau in TAU(i).
@ingroup magma_dgeqlf_comp
********************************************************************/
extern "C" magma_int_t
magma_dgeqlf(
magma_int_t m, magma_int_t n,
double *A, magma_int_t lda, double *tau,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) ( A + (i_) + (j_)*lda)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dwork(i_) (dwork + (i_))
magmaDouble_ptr dA, dwork;
double c_one = MAGMA_D_ONE;
magma_int_t i, k, lddwork, old_i, old_ib, nb;
magma_int_t rows, cols;
magma_int_t ib, ki, kk, mu, nu, iinfo, ldda;
int lquery;
nb = magma_get_dgeqlf_nb(m);
*info = 0;
lquery = (lwork == -1);
// silence "uninitialized" warnings
old_ib = nb;
old_i = 0;
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
}
k = min(m,n);
if (*info == 0) {
if (k == 0)
work[0] = c_one;
else {
work[0] = MAGMA_D_MAKE( max(n*nb, 2*nb*nb), 0 );
}
if (lwork < max(max(1,n), 2*nb*nb) && ! lquery)
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Quick return if possible */
if (k == 0)
return *info;
lddwork = ((n+31)/32)*32;
ldda = ((m+31)/32)*32;
if (MAGMA_SUCCESS != magma_dmalloc( &dA, (n)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dwork = dA + ldda*n;
magma_queue_t queues[2];
magma_queue_create( &queues[0] );
magma_queue_create( &queues[1] );
if ( (nb > 1) && (nb < k) ) {
/* Use blocked code initially.
The last kk columns are handled by the block method.
First, copy the matrix on the GPU except the last kk columns */
magma_dsetmatrix_async( m, n-nb,
A(0, 0), lda,
dA(0, 0), ldda, queues[0] );
ki = ((k - nb - 1) / nb) * nb;
kk = min(k, ki + nb);
for (i = k - kk + ki; i >= k -kk; i -= nb) {
ib = min(k-i,nb);
if (i < k - kk + ki) {
/* 1. Copy asynchronously the current panel to the CPU.
2. Copy asynchronously the submatrix below the panel
to the CPU) */
rows = m - k + i + ib;
magma_dgetmatrix_async( rows, ib,
dA(0, n-k+i), ldda,
A(0, n-k+i), lda, queues[1] );
magma_dgetmatrix_async( m-rows, ib,
dA(rows, n-k+i), ldda,
A(rows, n-k+i), lda, queues[0] );
/* Apply H' to A(1:m-k+i+ib-1,1:n-k+i-1) from the left in
two steps - implementing the lookahead techniques.
This is the main update from the lookahead techniques. */
rows = m - k + old_i + old_ib;
cols = n - k + old_i - old_ib;
magma_dlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, cols, old_ib,
dA(0, cols+old_ib), ldda, dwork(0), lddwork,
dA(0, 0 ), ldda, dwork(old_ib), lddwork);
}
magma_queue_sync( queues[1] );
/* Compute the QL factorization of the current block
A(1:m-k+i+ib-1,n-k+i:n-k+i+ib-1) */
rows = m - k + i + ib;
cols = n - k + i;
lapackf77_dgeqlf( &rows, &ib, A(0,cols), &lda, tau+i, work, &lwork, &iinfo );
if (cols > 0) {
/* Form the triangular factor of the block reflector
H = H(i+ib-1) . . . H(i+1) H(i) */
lapackf77_dlarft( MagmaBackwardStr, MagmaColumnwiseStr,
&rows, &ib,
A(0, cols), &lda, tau + i, work, &ib);
dpanel_to_q( MagmaLower, ib, A(rows-ib,cols), lda, work+ib*ib);
magma_dsetmatrix( rows, ib,
A(0,cols), lda,
dA(0,cols), ldda );
dq_to_panel( MagmaLower, ib, A(rows-ib,cols), lda, work+ib*ib);
// Send the triangular part on the GPU
magma_dsetmatrix( ib, ib, work, ib, dwork(0), lddwork );
/* Apply H' to A(1:m-k+i+ib-1,1:n-k+i-1) from the left in
two steps - implementing the lookahead techniques.
This is the update of first ib columns. */
if (i-ib >= k -kk)
magma_dlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, ib, ib,
dA(0, cols), ldda, dwork(0), lddwork,
dA(0,cols-ib), ldda, dwork(ib), lddwork);
else {
magma_dlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaBackward, MagmaColumnwise,
rows, cols, ib,
dA(0, cols), ldda, dwork(0), lddwork,
dA(0, 0 ), ldda, dwork(ib), lddwork);
}
old_i = i;
old_ib = ib;
}
}
mu = m - k + i + nb;
nu = n - k + i + nb;
magma_dgetmatrix( m, nu, dA(0,0), ldda, A(0,0), lda );
} else {
mu = m;
nu = n;
}
/* Use unblocked code to factor the last or only block */
if (mu > 0 && nu > 0)
lapackf77_dgeqlf(&mu, &nu, A(0,0), &lda, tau, work, &lwork, &iinfo);
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dA );
return *info;
} /* magma_dgeqlf */
/**
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 */
/**
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_sgeqrf2_gpu(
magma_int_t m, magma_int_t n,
magmaFloat_ptr dA, size_t dA_offset, magma_int_t ldda,
float *tau,
magma_queue_t* queue,
magma_int_t *info)
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
SGEQRF computes a QR factorization of a real M-by-N matrix A:
A = Q * R.
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.
dA (input/output) REAL array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
TAU (output) REAL array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define dA(a_1,a_2) dA, (dA_offset + (a_1) + (a_2)*(ldda))
#define work(a_1) ( work + (a_1))
#define hwork ( work + (nb)*(m))
magmaFloat_ptr dwork;
float *work;
magma_int_t i, k, ldwork, lddwork, old_i, old_ib, rows;
magma_int_t nbmin, nx, ib, nb;
magma_int_t lhwork, lwork;
*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;
}
k = min(m,n);
if (k == 0)
return *info;
nb = magma_get_sgeqrf_nb(m);
lwork = (m+n) * nb;
lhwork = lwork - (m)*nb;
if ( MAGMA_SUCCESS != magma_smalloc( &dwork, n*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/*
if ( MAGMA_SUCCESS != magma_smalloc_cpu( &work, lwork ) ) {
*info = MAGMA_ERR_HOST_ALLOC;
magma_free( dwork );
return *info;
}
*/
cl_mem buffer = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
sizeof(float)*lwork, NULL, NULL);
work = (float*)clEnqueueMapBuffer(queue[0], buffer, CL_TRUE,
CL_MAP_READ | CL_MAP_WRITE,
0, lwork*sizeof(float),
0, NULL, NULL, NULL);
nbmin = 2;
nx = nb;
ldwork = m;
lddwork= n;
if (nb >= nbmin && nb < k && nx < k) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nx; i += nb) {
ib = min(k-i, nb);
rows = m -i;
magma_queue_sync( queue[1] );
magma_sgetmatrix_async(rows, ib, dA(i, i), ldda, work(i), ldwork, queue[0], NULL);
if (i > 0) {
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
dA(old_i, old_i ), ldda, dwork,0, lddwork,
dA(old_i, old_i+2*old_ib), ldda, dwork,old_ib, lddwork, queue[1]);
magma_ssetmatrix_async( old_ib, old_ib, work(old_i), ldwork,
dA(old_i, old_i), ldda, queue[1], NULL);
}
magma_queue_sync(queue[0]);
lapackf77_sgeqrf(&rows, &ib, work(i), &ldwork, tau+i, hwork, &lhwork, info);
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
work(i), &ldwork, tau+i, hwork, &ib);
spanel_to_q( MagmaUpper, ib, work(i), ldwork, hwork+ib*ib );
/* download the i-th V matrix */
magma_ssetmatrix_async(rows, ib, work(i), ldwork, dA(i,i), ldda, queue[0], NULL);
/* download the T matrix */
magma_queue_sync( queue[1] );
magma_ssetmatrix_async( ib, ib, hwork, ib, dwork, 0, lddwork, queue[0], NULL);
magma_queue_sync( queue[0] );
if (i + ib < n) {
if (i+nb < k-nx) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left */
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
dA(i, i ), ldda, dwork,0, lddwork,
dA(i, i+ib), ldda, dwork,ib, lddwork, queue[1]);
sq_to_panel( MagmaUpper, ib, work(i), ldwork, hwork+ib*ib );
}
else {
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
dA(i, i ), ldda, dwork,0, lddwork,
dA(i, i+ib), ldda, dwork,ib, lddwork, queue[1]);
sq_to_panel( MagmaUpper, ib, work(i), ldwork, hwork+ib*ib );
magma_ssetmatrix_async(ib, ib, work(i), ldwork, dA(i,i), ldda, queue[1], NULL);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
magma_free(dwork);
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
magma_sgetmatrix_async(rows, ib, dA(i, i), ldda, work, rows, queue[1], NULL);
magma_queue_sync(queue[1]);
lhwork = lwork - rows*ib;
lapackf77_sgeqrf(&rows, &ib, work, &rows, tau+i, work+ib*rows, &lhwork, info);
magma_ssetmatrix_async(rows, ib, work, rows, dA(i, i), ldda, queue[1], NULL);
}
magma_queue_sync(queue[0]);
magma_queue_sync(queue[1]);
// magma_free_cpu(work);
clEnqueueUnmapMemObject(queue[0], buffer, work, 0, NULL, NULL);
clReleaseMemObject(buffer);
return *info;
} /* magma_sgeqrf2_gpu */
/***************************************************************************//**
Purpose
-------
DGEQRF computes a QR factorization of a real M-by-N matrix A:
A = Q * R.
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_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[in,out]
dR_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array on the GPU, dimension (LDDR, N/NB)
dR should be of size (LDDR, N) when provide_RT > 0 and
of size (LDDT, NB) otherwise. NB is the local blocking size.
On exit, the elements of R are stored in dR only when provide_RT > 0.
@param[in]
lddr INTEGER
The leading dimension of the array dR.
LDDR >= min(M,N) when provide_RT == 1
otherwise LDDR >= min(NB, min(M,N)).
NB is the local blocking size.
To benefit from coalescent memory accesses LDDR must be
divisible by 16.
@param[in,out]
dT_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array on the GPU, dimension (LDDT, N/NB)
dT should be of size (LDDT, N) when provide_RT > 0 and
of size (LDDT, NB) otherwise. NB is the local blocking size.
On exit, the elements of T are stored in dT only when provide_RT > 0.
@param[in]
lddt INTEGER
The leading dimension of the array dT.
LDDT >= min(NB,min(M,N)). NB is the local blocking size.
To benefit from coalescent memory accesses LDDR must be
divisible by 16.
@param[out]
dtau_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[in]
provide_RT INTEGER
provide_RT = 0 no R and no T in output.
dR and dT are used as local workspace to store the R and T of each step.
provide_RT = 1 the whole R of size (min(M,N), N) and the nbxnb block of T are provided in output.
provide_RT = 2 the nbxnb diag block of R and of T are provided in output.
@param[out]
info_array Array of INTEGERs, dimension (batchCount), for corresponding matrices.
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_geqrf_batched
*******************************************************************************/
extern "C" magma_int_t
magma_dgeqrf_expert_batched(
magma_int_t m, magma_int_t n,
double **dA_array, magma_int_t ldda,
double **dR_array, magma_int_t lddr,
double **dT_array, magma_int_t lddt,
double **dtau_array, magma_int_t provide_RT,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define dA(i, j) (dA + (i) + (j)*ldda)
/* Local Parameter */
magma_int_t nb = magma_get_dgeqrf_batched_nb(m);
magma_int_t nnb = 8;
magma_int_t min_mn = min(m, n);
/* Check arguments */
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
else if (lddr < min_mn && provide_RT == 1)
arginfo = -6;
else if (lddr < min(min_mn, nb))
arginfo = -6;
else if (lddt < min(min_mn, nb))
arginfo = -8;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if (min_mn == 0 ) return arginfo;
if ( m > 2048 || n > 2048 ) {
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t i, k, ib=nb, jb=nnb, offset_RT=0, use_stream;
magma_int_t ldw, offset;
double **dW0_displ = NULL;
double **dW1_displ = NULL;
double **dW2_displ = NULL;
double **dW3_displ = NULL;
double **dW4_displ = NULL;
double **dW5_displ = NULL;
double **dR_displ = NULL;
double **dT_displ = NULL;
double *dwork = NULL;
double **cpuAarray = NULL;
double **cpuTarray = NULL;
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dW5_displ, batchCount * sizeof(*dW5_displ));
magma_malloc((void**)&dR_displ, batchCount * sizeof(*dR_displ));
magma_malloc((void**)&dT_displ, batchCount * sizeof(*dT_displ));
magma_dmalloc(&dwork, (2 * nb * n) * batchCount);
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(double*));
magma_malloc_cpu((void**) &cpuTarray, batchCount*sizeof(double*));
/* check allocation */
if ( dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dW5_displ == NULL ||
dR_displ == NULL || dT_displ == NULL || dwork == NULL ||
cpuAarray == NULL || cpuTarray == NULL ) {
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magma_ddisplace_pointers(dR_displ, dR_array, lddr, 0, 0, batchCount, queue);
magma_ddisplace_pointers(dT_displ, dT_array, lddt, 0, 0, batchCount, queue);
// set dwork to zero because our GEMM routine does propagate NAN when C=betaC+alphaA*B and beta=0
magmablas_dlaset_q( MagmaFull, 2*nb, n*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dwork, 2*nb, queue );
// set dR and dT to zero. if provide_RT == 0 only a tile of size nbxnb is used and overwritten at each step
magmablas_dlaset_batched( MagmaFull, lddr, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_D_ZERO, MAGMA_D_ZERO, dR_displ, lddr, batchCount, queue );
magmablas_dlaset_batched( MagmaFull, lddt, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_D_ZERO, MAGMA_D_ZERO, dT_displ, lddt, batchCount, queue );
/*
if ( provide_RT > 0 )
{
magmablas_dlaset_q( MagmaFull, lddr, n*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dR, lddr, queue );
magmablas_dlaset_q( MagmaFull, lddt, n*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dT, lddt, queue );
}
else
{
magmablas_dlaset_q( MagmaFull, lddr, nb*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dR, lddr, queue );
magmablas_dlaset_q( MagmaFull, lddt, nb*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dT, lddt, queue );
}
*/
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (i=0; i < nbstreams; i++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[i] );
}
magma_getvector( batchCount, sizeof(double*), dA_array, 1, cpuAarray, 1, queue);
magma_getvector( batchCount, sizeof(double*), dT_array, 1, cpuTarray, 1, queue);
for (i=0; i < min_mn; i += nb)
{
ib = min(nb, min_mn-i);
//===============================================
// panel factorization
//===============================================
magma_ddisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
magma_ddisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
if ( provide_RT > 0 )
{
offset_RT = i;
magma_ddisplace_pointers(dR_displ, dR_array, lddr, (provide_RT == 1 ? offset_RT:0), offset_RT, batchCount, queue);
magma_ddisplace_pointers(dT_displ, dT_array, lddt, 0, offset_RT, batchCount, queue);
}
//dwork is used in panel factorization and trailing matrix update
//dW4_displ, dW5_displ are used as workspace and configured inside
magma_dgeqrf_panel_batched(m-i, ib, jb,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dR_displ, lddr,
dW1_displ,
dW3_displ,
dwork,
dW4_displ, dW5_displ,
info_array,
batchCount, queue);
//===============================================
// end of panel
//===============================================
//===============================================
// update trailing matrix
//===============================================
if ( (n-ib-i) > 0)
{
//dwork is used in panel factorization and trailing matrix update
//reset dW4_displ
ldw = nb;
magma_dset_pointer( dW4_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
offset = ldw*n*batchCount;
magma_dset_pointer( dW5_displ, dwork + offset, 1, 0, 0, ldw*n, batchCount, queue );
// set the diagonal of v as one and the upper triangular part as zero already set inside geqrf_panel
//magmablas_dlaset_batched( MagmaUpper, ib, ib, MAGMA_D_ZERO, MAGMA_D_ONE, dW0_displ, ldda, batchCount, queue );
//magma_ddisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
// it is faster since it is using BLAS-3 GEMM routines, different from lapack implementation
magma_dlarft_batched(m-i, ib, 0,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dW4_displ, nb*lddt,
batchCount, queue);
// perform C = (I-V T^H V^H) * C, C is the trailing matrix
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
use_stream = magma_drecommend_cublas_gemm_stream(MagmaNoTrans, MagmaNoTrans, m-i-ib, n-i-ib, ib);
if ( use_stream )
{
magma_queue_sync(queue);
for (k=0; k < batchCount; k++)
{
streamid = k%nbstreams;
// the queue gemm must take cpu pointer
magma_dlarfb_gpu_gemm( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
cpuAarray[k] + i + i * ldda, ldda,
cpuTarray[k] + offset_RT*lddt, lddt,
cpuAarray[k] + i + (i+ib) * ldda, ldda,
dwork + nb * n * k, -1,
dwork + nb * n * batchCount + nb * n * k, -1, queues[streamid] );
}
// need to synchronise to be sure that panel does not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
//direct trailing matrix in dW1_displ
magma_ddisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_dlarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
(const double**)dW0_displ, ldda,
(const double**)dT_displ, lddt,
dW1_displ, ldda,
dW4_displ, ldw,
dW5_displ, ldw,
batchCount, queue );
}
}// update the trailing matrix
//===============================================
// copy dR back to V after the trailing matrix update,
// only when provide_RT=0 otherwise the nbxnb block of V is set to diag=1/0
// The upper portion of V could be set totaly to 0 here
if ( provide_RT == 0 )
{
magmablas_dlacpy_batched( MagmaUpper, ib, ib, dR_displ, lddr, dW0_displ, ldda, batchCount, queue );
}
}
magma_queue_sync(queue);
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
return arginfo;
}
/**
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_cgeqrf_ooc(magma_int_t m, magma_int_t n,
magmaFloatComplex *a, magma_int_t lda, magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magma_int_t *info )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGEQRF_OOC computes a QR factorization of a COMPLEX M-by-N matrix A:
A = Q * R. This version does not require work space on the GPU
passed as input. GPU memory is allocated in the routine.
This is an out-of-core (ooc) version that is similar to magma_cgeqrf but
the difference is that this version can use a GPU even if the matrix
does not fit into the GPU memory at once.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
Higher performance is achieved if WORK is in pinned memory, e.g.
allocated using magma_malloc_pinned.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_cgeqrf_nb(M).
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a+(a_2)*(lda) + (a_1))
#define da_ref(a_1,a_2) (da+(a_2)*ldda + (a_1))
magmaFloatComplex *da, *dwork;
magmaFloatComplex c_one = MAGMA_C_ONE;
int k, lddwork, ldda;
*info = 0;
int nb = magma_get_cgeqrf_nb(min(m, n));
int lwkopt = n * nb;
work[0] = MAGMA_C_MAKE( (float)lwkopt, 0 );
int lquery = (lwork == -1);
if (m < 0) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,m)) {
*info = -4;
} else if (lwork < max(1,n) && ! lquery) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
/* Check how much memory do we have */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
magma_int_t IB, NB = (magma_int_t)(0.8*freeMem/m);
NB = (NB / nb) * nb;
if (NB >= n)
return magma_cgeqrf(m, n, a, lda, tau, work, lwork, info);
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return *info;
}
lddwork = ((NB+31)/32)*32+nb;
ldda = ((m+31)/32)*32;
if (MAGMA_SUCCESS != magma_cmalloc( &da, (NB + nb)*ldda + nb*lddwork )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
// magmablasSetKernelStream(stream[1]);
magmaFloatComplex *ptr = da + ldda * NB;
dwork = da + ldda*(NB + nb);
/* start the main loop over the blocks that fit in the GPU memory */
for(int i=0; i<n; i+=NB) {
IB = min(n-i, NB);
//printf("Processing %5d columns -- %5d to %5d ... \n", IB, i, i+IB);
/* 1. Copy the next part of the matrix to the GPU */
magma_csetmatrix_async( (m), IB,
a_ref(0,i), lda,
da_ref(0,0), ldda, stream[0] );
magma_queue_sync( stream[0] );
/* 2. Update it with the previous transformations */
for(int j=0; j<min(i,k); j+=nb) {
magma_int_t ib = min(k-j, nb);
/* Get a panel in ptr. */
// 1. Form the triangular factor of the block reflector
// 2. Send it to the GPU.
// 3. Put 0s in the upper triangular part of V.
// 4. Send V to the GPU in ptr.
// 5. Update the matrix.
// 6. Restore the upper part of V.
magma_int_t rows = m-j;
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, a_ref(j,j), &lda, tau+j, work, &ib);
magma_csetmatrix_async( ib, ib,
work, ib,
dwork, lddwork, stream[1] );
cpanel_to_q(MagmaUpper, ib, a_ref(j,j), lda, work+ib*ib);
magma_csetmatrix_async( rows, ib,
a_ref(j,j), lda,
ptr, rows, stream[1] );
magma_queue_sync( stream[1] );
magma_clarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, IB, ib,
ptr, rows, dwork, lddwork,
da_ref(j, 0), ldda, dwork+ib, lddwork);
cq_to_panel(MagmaUpper, ib, a_ref(j,j), lda, work+ib*ib);
}
/* 3. Do a QR on the current part */
if (i<k)
magma_cgeqrf2_gpu(m-i, IB, da_ref(i,0), ldda, tau+i, info);
/* 4. Copy the current part back to the CPU */
magma_cgetmatrix_async( (m), IB,
da_ref(0,0), ldda,
a_ref(0,i), lda, stream[0] );
}
magma_queue_sync( stream[0] );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( da );
return *info;
} /* magma_cgeqrf_ooc */
extern "C" magma_int_t
magma_ssygvd(magma_int_t itype, char jobz, char uplo, magma_int_t n,
float *a, magma_int_t lda, float *b, magma_int_t ldb,
float *w, float *work, magma_int_t lwork,
magma_int_t *iwork, magma_int_t liwork, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SSYGVD computes all the eigenvalues, and optionally, the eigenvectors
of a real generalized symmetric-definite eigenproblem, of the form
A*x=(lambda)*B*x, A*Bx=(lambda)*x, or B*A*x=(lambda)*x. Here A and
B are assumed to be symmetric and B is also positive definite.
If eigenvectors are desired, it uses a divide and conquer algorithm.
The divide and conquer algorithm 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
=========
ITYPE (input) INTEGER
Specifies the problem type to be solved:
= 1: A*x = (lambda)*B*x
= 2: A*B*x = (lambda)*x
= 3: B*A*x = (lambda)*x
JOBZ (input) CHARACTER*1
= 'N': Compute eigenvalues only;
= 'V': Compute eigenvalues and eigenvectors.
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A and B are stored;
= 'L': Lower triangles of A and B are stored.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
A (input/output) COMPLEX*16 array, dimension (LDA, N)
On entry, the 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. If UPLO = 'L',
the leading N-by-N lower triangular part of A contains
the lower triangular part of the matrix A.
On exit, if JOBZ = 'V', then if INFO = 0, A contains the
matrix Z of eigenvectors. The eigenvectors are normalized
as follows:
if ITYPE = 1 or 2, Z**T * B * Z = I;
if ITYPE = 3, Z**T * inv(B) * Z = I.
If JOBZ = 'N', then on exit the upper triangle (if UPLO='U')
or the lower triangle (if UPLO='L') of A, including the
diagonal, is destroyed.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input/output) COMPLEX*16 array, dimension (LDB, N)
On entry, the symmetric matrix B. If UPLO = 'U', the
leading N-by-N upper triangular part of B contains the
upper triangular part of the matrix B. If UPLO = 'L',
the leading N-by-N lower triangular part of B contains
the lower triangular part of the matrix B.
On exit, if INFO <= N, the part of B containing the matrix is
overwritten by the triangular factor U or L from the Cholesky
factorization B = U**T * U or B = L * L**T.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
W (output) DOUBLE PRECISION array, dimension (N)
If INFO = 0, the eigenvalues in ascending order.
WORK (workspace/output) COMPLEX*16 array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The length of the array WORK.
If N <= 1, LWORK >= 1.
If JOBZ = 'N' and N > 1, LWORK >= 2*N*nb + 1.
If JOBZ = 'V' and N > 1, LWORK >= 1 + 6*N*nb + 2*N**2.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal sizes of the WORK and
IWORK arrays, returns these values as the first entries of
the WORK and IWORK arrays, and no error message
related to LWORK or LIWORK is issued by XERBLA.
IWORK (workspace/output) INTEGER array, dimension (MAX(1,LIWORK))
On exit, if INFO = 0, IWORK(1) returns the optimal LIWORK.
LIWORK (input) INTEGER
The dimension of the array IWORK.
If N <= 1, LIWORK >= 1.
If JOBZ = 'N' and N > 1, LIWORK >= 1.
If JOBZ = 'V' and N > 1, LIWORK >= 3 + 5*N.
If LIWORK = -1, then a workspace query is assumed; the
routine only calculates the optimal sizes of the WORK
and IWORK arrays, returns these values as the first entries
of the WORK and IWORK arrays, and no error message
related to LWORK or LIWORK is issued by XERBLA.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: SPOTRF or SSYEVD returned an error code:
<= N: if INFO = i and JOBZ = 'N', then the algorithm
failed to converge; i off-diagonal elements of an
intermediate tridiagonal form did not converge to
zero;
if INFO = i and JOBZ = 'V', then the algorithm
failed to compute an eigenvalue while working on
the submatrix lying in rows and columns INFO/(N+1)
through mod(INFO,N+1);
> N: if INFO = N + i, for 1 <= i <= N, then the leading
minor of order i of B is not positive definite.
The factorization of B could not be completed and
no eigenvalues or eigenvectors were computed.
Further Details
===============
Based on contributions by
Mark Fahey, Department of Mathematics, Univ. of Kentucky, USA
Modified so that no backsubstitution is performed if SSYEVD fails to
converge (NEIG in old code could be greater than N causing out of
bounds reference to A - reported by Ralf Meyer). Also corrected the
description of INFO and the test on ITYPE. Sven, 16 Feb 05.
===================================================================== */
char uplo_[2] = {uplo, 0};
char jobz_[2] = {jobz, 0};
float d_one = MAGMA_S_ONE;
float *da;
float *db;
magma_int_t ldda = n;
magma_int_t lddb = n;
magma_int_t lower;
char trans[1];
magma_int_t wantz, lquery;
magma_int_t lopt, lwmin, liopt, liwmin;
cudaStream_t stream;
magma_queue_create( &stream );
wantz = lapackf77_lsame(jobz_, MagmaVectorsStr);
lower = lapackf77_lsame(uplo_, MagmaLowerStr);
lquery = lwork == -1 || liwork == -1;
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! (wantz || lapackf77_lsame(jobz_, MagmaNoVectorsStr))) {
*info = -2;
} else if (! (lower || lapackf77_lsame(uplo_, MagmaUpperStr))) {
*info = -3;
} else if (n < 0) {
*info = -4;
} else if (lda < max(1,n)) {
*info = -6;
} else if (ldb < max(1,n)) {
*info = -8;
}
magma_int_t nb = magma_get_ssytrd_nb(n);
if (n < 1) {
liwmin = 1;
lwmin = 1;
} else if (wantz) {
lwmin = 1 + 6 * n * nb + 2* n * n;
liwmin = 5 * n + 3;
} else {
lwmin = 2 * n * nb + 1;
liwmin = 1;
}
lopt = lwmin;
liopt = liwmin;
work[ 0] = lopt;
iwork[0] = liopt;
if (lwork < lwmin && ! lquery) {
*info = -11;
} else if (liwork < liwmin && ! lquery) {
*info = -13;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return MAGMA_ERR_ILLEGAL_VALUE;
}
else if (lquery) {
return MAGMA_SUCCESS;
}
/* Quick return if possible */
if (n == 0) {
return 0;
}
if (MAGMA_SUCCESS != magma_smalloc( &da, n*ldda ) ||
MAGMA_SUCCESS != magma_smalloc( &db, n*lddb )) {
*info = -17;
return MAGMA_ERR_DEVICE_ALLOC;
}
/* Form a Cholesky factorization of B. */
magma_ssetmatrix( n, n, b, ldb, db, lddb );
magma_ssetmatrix_async( n, n,
a, lda,
da, ldda, stream );
magma_spotrf_gpu(uplo_[0], n, db, lddb, info);
if (*info != 0) {
*info = n + *info;
return 0;
}
magma_queue_sync( stream );
magma_sgetmatrix_async( n, n,
db, lddb,
b, ldb, stream );
/* Transform problem to standard eigenvalue problem and solve. */
magma_ssygst_gpu(itype, uplo_[0], n, da, ldda, db, lddb, info);
magma_ssyevd_gpu(jobz_[0], uplo_[0], n, da, ldda, w, a, lda,
work, lwork, iwork, liwork, info);
lopt = max( lopt, (magma_int_t) work[0]);
liopt = max(liopt, iwork[0]);
if (wantz && *info == 0)
{
/* Backtransform eigenvectors to the original problem. */
if (itype == 1 || itype == 2)
{
/* For A*x=(lambda)*B*x and A*B*x=(lambda)*x;
backtransform eigenvectors: x = inv(L)'*y or inv(U)*y */
if (lower) {
*(unsigned char *)trans = MagmaTrans;
} else {
*(unsigned char *)trans = MagmaNoTrans;
}
magma_strsm(MagmaLeft, uplo_[0], *trans, MagmaNonUnit,
n, n, d_one, db, lddb, da, ldda);
} else if (itype == 3)
{
/* For B*A*x=(lambda)*x;
backtransform eigenvectors: x = L*y or U'*y */
if (lower) {
*(unsigned char *)trans = MagmaNoTrans;
} else {
*(unsigned char *)trans = MagmaTrans;
}
magma_strmm(MagmaLeft, uplo_[0], *trans, MagmaNonUnit,
n, n, d_one, db, lddb, da, ldda);
}
magma_sgetmatrix( n, n, da, ldda, a, lda );
}
magma_queue_sync( stream );
magma_queue_destroy( stream );
work[0] = (float) lopt;
iwork[0] = liopt;
magma_free( da );
magma_free( db );
return MAGMA_SUCCESS;
} /* magma_ssygvd */
/**
Purpose
-------
CLATRD2 reduces NB rows and columns of a complex Hermitian matrix A to
Hermitian tridiagonal form by an orthogonal similarity
transformation Q' * A * Q, and returns the matrices V and W which are
needed to apply the transformation to the unreduced part of A.
If UPLO = MagmaUpper, CLATRD reduces the last NB rows and columns of a
matrix, of which the upper triangle is supplied;
if UPLO = MagmaLower, CLATRD reduces the first NB rows and columns of a
matrix, of which the lower triangle is supplied.
This is an auxiliary routine called by CHETRD2_GPU. It uses an
accelerated HEMV that needs extra memory.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the upper or lower triangular part of the
Hermitian matrix A is stored:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the matrix A.
@param[in]
nb INTEGER
The number of rows and columns to be reduced.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
n-by-n upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading n-by-n lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit:
- if UPLO = MagmaUpper, the last NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements above the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors;
- if UPLO = MagmaLower, the first NB columns have been reduced to
tridiagonal form, with the diagonal elements overwriting
the diagonal elements of A; the elements below the diagonal
with the array TAU, represent the orthogonal matrix Q as a
product of elementary reflectors.
See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= (1,N).
@param[out]
e COMPLEX array, dimension (N-1)
If UPLO = MagmaUpper, E(n-nb:n-1) contains the superdiagonal
elements of the last NB columns of the reduced matrix;
if UPLO = MagmaLower, E(1:nb) contains the subdiagonal elements of
the first NB columns of the reduced matrix.
@param[out]
tau COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors, stored in
TAU(n-nb:n-1) if UPLO = MagmaUpper, and in TAU(1:nb) if UPLO = MagmaLower.
See Further Details.
@param[out]
W COMPLEX array, dimension (LDW,NB)
The n-by-nb matrix W required to update the unreduced part
of A.
@param[in]
ldw INTEGER
The leading dimension of the array W. LDW >= max(1,N).
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n) H(n-1) . . . H(n-nb+1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i:n) = 0 and v(i-1) = 1; v(1:i-1) is stored on exit in A(1:i-1,i),
and tau in TAU(i-1).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(nb).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+1:n) is stored on exit in A(i+1:n,i),
and tau in TAU(i).
The elements of the vectors v together form the n-by-nb matrix V
which is needed, with W, to apply the transformation to the unreduced
part of the matrix, using a Hermitian rank-2k update of the form:
A := A - V*W' - W*V'.
The contents of A on exit are illustrated by the following examples
with n = 5 and nb = 2:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( a a a v4 v5 ) ( d )
( a a v4 v5 ) ( 1 d )
( a 1 v5 ) ( v1 1 a )
( d 1 ) ( v1 v2 a a )
( d ) ( v1 v2 a a a )
where d denotes a diagonal element of the reduced matrix, a denotes
an element of the original matrix that is unchanged, and vi denotes
an element of the vector defining H(i).
@ingroup magma_cheev_aux
********************************************************************/
extern "C" magma_int_t
magma_clatrd2(magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
float *e, magmaFloatComplex *tau,
magmaFloatComplex *W, magma_int_t ldw,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dW, magma_int_t lddw,
magmaFloatComplex *dwork, magma_int_t ldwork)
{
#define A(i, j) (A + (j)*lda + (i))
#define W(i, j) (W + (j)*ldw + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
#define dW(i, j) (dW + (j)*lddw + (i))
magma_int_t i;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex value = MAGMA_C_ZERO;
magma_int_t ione = 1;
magma_int_t i_n, i_1, iw;
magmaFloatComplex alpha;
magmaFloatComplex *f;
if (n <= 0) {
return 0;
}
magma_queue_t stream;
magma_queue_create( &stream );
magma_cmalloc_cpu( &f, n );
assert( f != NULL ); // TODO return error, or allocate outside clatrd
if (uplo == MagmaUpper) {
/* Reduce last NB columns of upper triangle */
for (i = n-1; i >= n - nb; --i) {
i_1 = i + 1;
i_n = n - i - 1;
iw = i - n + nb;
if (i < n-1) {
/* Update A(1:i,i) */
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, W(i, iw+1), &ldw);
#endif
blasf77_cgemv("No transpose", &i_1, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i, iw+1), &ldw, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, W(i, iw+1), &ldw);
lapackf77_clacgv(&i_n, A(i, i+1), &ldw);
#endif
blasf77_cgemv("No transpose", &i_1, &i_n, &c_neg_one, W(0, iw+1), &ldw,
A(i, i+1), &lda, &c_one, A(0, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i_n, A(i, i+1), &ldw);
#endif
}
if (i > 0) {
/* Generate elementary reflector H(i) to annihilate A(1:i-2,i) */
alpha = *A(i-1, i);
lapackf77_clarfg(&i, &alpha, A(0, i), &ione, &tau[i - 1]);
e[i-1] = MAGMA_C_REAL( alpha );
*A(i-1,i) = MAGMA_C_MAKE( 1, 0 );
/* Compute W(1:i-1,i) */
// 1. Send the block reflector A(0:n-i-1,i) to the GPU
magma_csetvector( i, A(0, i), 1, dA(0, i), 1 );
//#if (GPUSHMEM < 200)
//magma_chemv(MagmaUpper, i, c_one, dA(0, 0), ldda,
// dA(0, i), ione, c_zero, dW(0, iw), ione);
//#else
magmablas_chemv_work(MagmaUpper, i, c_one, dA(0, 0), ldda,
dA(0, i), ione, c_zero, dW(0, iw), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_cgetmatrix_async( i, 1,
dW(0, iw), lddw,
W(0, iw) /*test*/, ldw, stream );
if (i < n-1) {
blasf77_cgemv(MagmaConjTransStr, &i, &i_n, &c_one, W(0, iw+1), &ldw,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
}
// 3. Here is where we need it // TODO find the right place
magma_queue_sync( stream );
if (i < n-1) {
blasf77_cgemv("No transpose", &i, &i_n, &c_neg_one, A(0, i+1), &lda,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
blasf77_cgemv(MagmaConjTransStr, &i, &i_n, &c_one, A(0, i+1), &lda,
A(0, i), &ione, &c_zero, W(i+1, iw), &ione);
blasf77_cgemv("No transpose", &i, &i_n, &c_neg_one, W(0, iw+1), &ldw,
W(i+1, iw), &ione, &c_one, W(0, iw), &ione);
}
blasf77_cscal(&i, &tau[i - 1], W(0, iw), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i, W(0,iw), ione, A(0,i), ione, &value );
#else
value = cblas_cdotc( i, W(0,iw), ione, A(0,i), ione );
#endif
alpha = tau[i - 1] * -0.5f * value;
blasf77_caxpy(&i, &alpha, A(0, i), &ione,
W(0, iw), &ione);
}
}
}
else {
/* Reduce first NB columns of lower triangle */
for (i = 0; i < nb; ++i) {
/* Update A(i:n,i) */
i_n = n - i;
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, W(i, 0), &ldw);
#endif
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, A(i, 0), &lda,
W(i, 0), &ldw, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, W(i, 0), &ldw);
lapackf77_clacgv(&i, A(i, 0), &lda);
#endif
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, W(i, 0), &ldw,
A(i, 0), &lda, &c_one, A(i, i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
lapackf77_clacgv(&i, A(i, 0), &lda);
#endif
if (i < n-1) {
/* Generate elementary reflector H(i) to annihilate A(i+2:n,i) */
i_n = n - i - 1;
alpha = *A(i+1, i);
lapackf77_clarfg(&i_n, &alpha, A(min(i+2,n-1), i), &ione, &tau[i]);
e[i] = MAGMA_C_REAL( alpha );
*A(i+1,i) = MAGMA_C_MAKE( 1, 0 );
/* Compute W(i+1:n,i) */
// 1. Send the block reflector A(i+1:n,i) to the GPU
magma_csetvector( i_n, A(i+1, i), 1, dA(i+1, i), 1 );
//#if (GPUSHMEM < 200)
//magma_chemv(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
// dW(i+1, i), ione);
//#else
magmablas_chemv_work(MagmaLower, i_n, c_one, dA(i+1, i+1), ldda, dA(i+1, i), ione, c_zero,
dW(i+1, i), ione,
dwork, ldwork);
//#endif
// 2. Start putting the result back (asynchronously)
magma_cgetmatrix_async( i_n, 1,
dW(i+1, i), lddw,
W(i+1, i), ldw, stream );
blasf77_cgemv(MagmaConjTransStr, &i_n, &i, &c_one, W(i+1, 0), &ldw,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, A(i+1, 0), &lda,
W(0, i), &ione, &c_zero, f, &ione);
blasf77_cgemv(MagmaConjTransStr, &i_n, &i, &c_one, A(i+1, 0), &lda,
A(i+1, i), &ione, &c_zero, W(0, i), &ione);
// 3. Here is where we need it
magma_queue_sync( stream );
if (i != 0)
blasf77_caxpy(&i_n, &c_one, f, &ione, W(i+1, i), &ione);
blasf77_cgemv("No transpose", &i_n, &i, &c_neg_one, W(i+1, 0), &ldw,
W(0, i), &ione, &c_one, W(i+1, i), &ione);
blasf77_cscal(&i_n, &tau[i], W(i+1,i), &ione);
#if defined(PRECISION_z) || defined(PRECISION_c)
cblas_cdotc_sub( i_n, W(i+1,i), ione, A(i+1,i), ione, &value );
#else
value = cblas_cdotc( i_n, W(i+1,i), ione, A(i+1,i), ione );
#endif
alpha = tau[i] * -0.5f * value;
blasf77_caxpy(&i_n, &alpha, A(i+1, i), &ione, W(i+1,i), &ione);
}
}
}
magma_free_cpu(f);
magma_queue_destroy( stream );
return 0;
} /* magma_clatrd */
/**
Purpose
-------
SGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
It uses 2 queues to overlap communication and computation.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A REAL array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_sgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_sgetrf(
magma_int_t m, magma_int_t n,
float *A, magma_int_t lda,
magma_int_t *ipiv,
magma_int_t *info)
{
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_)*nb + (j_)*nb*ldda + dA_offset)
#define dAT(i_, j_) dAT, ((i_)*nb*lddat + (j_)*nb + dAT_offset)
#define dwork(i_) dwork, (i_)
#else
#define dA(i_, j_) ( dA + (i_)*nb + (j_)*nb*ldda)
#define dAT(i_, j_) ( dAT + (i_)*nb*lddat + (j_)*nb)
#define dwork(i_) (dwork + (i_))
#endif
// Constants
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
// Local variables
float *work;
magmaFloat_ptr dA, dAT, dwork;
magma_int_t iinfo, nb;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
nb = magma_get_sgetrf_nb( m, n );
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_sgetrf( &m, &n, A, &lda, ipiv, info );
}
else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, lddat, maxdim;
magma_int_t i, j, rows, cols, s = min(m, n)/nb;
maxm = magma_roundup( m, 32 );
maxn = magma_roundup( n, 32 );
maxdim = max( maxm, maxn );
lddat = maxn;
ldda = maxm;
/* set number of GPUs */
magma_int_t ngpu = magma_num_gpus();
if ( ngpu > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
magma_queue_t queues[2] = { NULL, NULL };
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
/* check the memory requirement */
size_t mem_size = magma_queue_mem_size( queues[0] );
mem_size /= sizeof(float);
magma_int_t h = 1+(2+ngpu);
magma_int_t ngpu2 = ngpu;
magma_int_t NB = (magma_int_t)(0.8*mem_size/maxm - h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
if ( ngpu > ceil((float)NB/nb) ) {
ngpu2 = (magma_int_t)ceil((float)NB/nb);
h = 1+(2+ngpu2);
NB = (magma_int_t)(0.8*mem_size/maxm - h*nb);
}
if ( ngpu2*NB < n ) {
/* require too much memory, so call non-GPU-resident version */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
work = A;
if (maxdim*maxdim < 2*maxm*maxn) {
// if close to square, allocate square matrix and transpose in-place
// dwork is nb*maxm for panel, and maxdim*maxdim for A
if (MAGMA_SUCCESS != magma_smalloc( &dwork, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
dA = dwork + nb*maxm;
ldda = lddat = maxdim;
magma_ssetmatrix( m, n, A, lda, dA(0,0), ldda, queues[0] );
dAT = dA;
magmablas_stranspose_inplace( maxdim, dAT(0,0), lddat, queues[0] );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
// dwork is nb*maxm for panel, and maxm*maxn for A
if (MAGMA_SUCCESS != magma_smalloc( &dwork, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
dA = dwork + nb*maxm;
magma_ssetmatrix( m, n, A, lda, dA(0,0), ldda, queues[0] );
if (MAGMA_SUCCESS != magma_smalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dwork );
magma_sgetrf_m( ngpu, m, n, A, lda, ipiv, info );
return *info;
}
magmablas_stranspose( m, n, dA(0,0), ldda, dAT(0,0), lddat, queues[0] );
}
lapackf77_sgetrf( &m, &nb, work, &lda, ipiv, &iinfo );
for( j = 0; j < s; j++ ) {
// get j-th panel from device
cols = maxm - j*nb;
if (j > 0) {
magmablas_stranspose( nb, cols, dAT(j,j), lddat, dwork(0), cols, queues[0] );
magma_queue_sync( queues[0] );
magma_sgetmatrix_async( m-j*nb, nb, dwork(0), cols, work, lda, queues[1] );
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (j+1)*nb, nb,
c_one, dAT(j-1,j-1), lddat,
dAT(j-1,j+1), lddat, queues[0] );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-j*nb, nb,
c_neg_one, dAT(j-1,j+1), lddat,
dAT(j, j-1), lddat,
c_one, dAT(j, j+1), lddat, queues[0] );
// do the cpu part
rows = m - j*nb;
magma_queue_sync( queues[1] );
lapackf77_sgetrf( &rows, &nb, work, &lda, ipiv+j*nb, &iinfo );
}
if (*info == 0 && iinfo > 0)
*info = iinfo + j*nb;
// put j-th panel onto device
magma_ssetmatrix_async( m-j*nb, nb, work, lda, dwork(0), cols, queues[1] );
for( i=j*nb; i < j*nb + nb; ++i ) {
ipiv[i] += j*nb;
}
magmablas_slaswp( n, dAT(0,0), lddat, j*nb + 1, j*nb + nb, ipiv, 1, queues[0] );
magma_queue_sync( queues[1] );
magmablas_stranspose( cols, nb, dwork(0), cols, dAT(j,j), lddat, queues[0] );
// do the small non-parallel computations (next panel update)
if (s > (j+1)) {
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat, queues[0] );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat, queues[0] );
}
else {
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(j, j ), lddat,
dAT(j, j+1), lddat, queues[0] );
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
n-(j+1)*nb, m-(j+1)*nb, nb,
c_neg_one, dAT(j, j+1), lddat,
dAT(j+1, j ), lddat,
c_one, dAT(j+1, j+1), lddat, queues[0] );
}
}
magma_int_t nb0 = min( m - s*nb, n - s*nb );
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_stranspose( nb0, rows, dAT(s,s), lddat, dwork(0), cols, queues[0] );
magma_sgetmatrix_async( rows, nb0, dwork(0), cols, work, lda, queues[0] );
magma_queue_sync( queues[0] );
// do the cpu part
lapackf77_sgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo );
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
for( i=s*nb; i < s*nb + nb0; ++i ) {
ipiv[i] += s*nb;
}
magmablas_slaswp( n, dAT(0,0), lddat, s*nb + 1, s*nb + nb0, ipiv, 1, queues[0] );
// put j-th panel onto device
magma_ssetmatrix_async( rows, nb0, work, lda, dwork(0), cols, queues[0] );
magmablas_stranspose( rows, nb0, dwork(0), cols, dAT(s,s), lddat, queues[0] );
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s, s), lddat,
dAT(s, s)+nb0, lddat, queues[0] );
}
// undo transpose
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_stranspose_inplace( maxdim, dAT(0,0), lddat, queues[0] );
magma_sgetmatrix( m, n, dAT(0,0), lddat, A, lda, queues[0] );
}
else {
magmablas_stranspose( n, m, dAT(0,0), lddat, dA(0,0), ldda, queues[0] );
magma_sgetmatrix( m, n, dA(0,0), ldda, A, lda, queues[0] );
magma_free( dAT );
}
magma_free( dwork );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
}
return *info;
} /* magma_sgetrf */
/**
Purpose
-------
ZPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_zposv_comp
********************************************************************/
extern "C" magma_int_t
magma_zpotrf_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex_ptr dA, magma_int_t ldda,
magma_int_t *info )
{
#ifdef HAVE_clBLAS
#define dA(i_, j_) dA, ((i_) + (j_)*ldda + dA_offset)
#else
#define dA(i_, j_) (dA + (i_) + (j_)*ldda)
#endif
/* Constants */
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
const double d_one = 1.0;
const double d_neg_one = -1.0;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
bool upper = (uplo == MagmaUpper);
magma_int_t j, jb, nb;
magmaDoubleComplex *work;
*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 *info;
}
nb = magma_get_zpotrf_nb( n );
if (MAGMA_SUCCESS != magma_zmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
if (nb <= 1 || nb >= n) {
/* Use unblocked code. */
magma_zgetmatrix( n, n, dA(0,0), ldda, work, n, queues[0] );
lapackf77_zpotrf( uplo_, &n, work, &n, info );
magma_zsetmatrix( n, n, work, n, dA(0,0), ldda, queues[0] );
}
else {
/* Use blocked code. */
if (upper) {
//=========================================================
/* Compute the Cholesky factorization A = U'*U. */
for (j=0; j < n; j += nb) {
// apply all previous updates to diagonal block,
// then transfer it to CPU
jb = min( nb, n-j );
magma_zherk( MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dA(0, j), ldda,
d_one, dA(j, j), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[0] );
// apply all previous updates to block row right of diagonal block
if (j+jb < n) {
magma_zgemm( MagmaConjTrans, MagmaNoTrans,
jb, n-j-jb, j,
c_neg_one, dA(0, j ), ldda,
dA(0, j+jb), ldda,
c_one, dA(j, j+jb), ldda, queues[1] );
}
// simultaneous with above zgemm, transfer diagonal block,
// factor it on CPU, and test for positive definiteness
magma_queue_sync( queues[0] );
lapackf77_zpotrf( MagmaUpperStr, &jb, work, &jb, info );
magma_zsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, queues[1] );
if (*info != 0) {
*info = *info + j;
break;
}
// apply diagonal block to block row right of diagonal block
if (j+jb < n) {
magma_ztrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, n-j-jb,
c_one, dA(j, j), ldda,
dA(j, j+jb), ldda, queues[1] );
}
}
}
else {
//=========================================================
// Compute the Cholesky factorization A = L*L'.
for (j=0; j < n; j += nb) {
// apply all previous updates to diagonal block,
// then transfer it to CPU
jb = min( nb, n-j );
magma_zherk( MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dA(j, 0), ldda,
d_one, dA(j, j), ldda, queues[1] );
magma_queue_sync( queues[1] );
magma_zgetmatrix_async( jb, jb,
dA(j, j), ldda,
work, jb, queues[0] );
// apply all previous updates to block column below diagonal block
if (j+jb < n) {
magma_zgemm( MagmaNoTrans, MagmaConjTrans,
n-j-jb, jb, j,
c_neg_one, dA(j+jb, 0), ldda,
dA(j, 0), ldda,
c_one, dA(j+jb, j), ldda, queues[1] );
}
// simultaneous with above zgemm, transfer diagonal block,
// factor it on CPU, and test for positive definiteness
magma_queue_sync( queues[0] );
lapackf77_zpotrf( MagmaLowerStr, &jb, work, &jb, info );
magma_zsetmatrix_async( jb, jb,
work, jb,
dA(j, j), ldda, queues[1] );
if (*info != 0) {
*info = *info + j;
break;
}
// apply diagonal block to block column below diagonal
if (j+jb < n) {
magma_ztrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-j-jb, jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda, queues[1] );
}
}
}
}
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free_pinned( work );
return *info;
} /* magma_zpotrf_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 */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf_mgpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double error, work[1];
double c_neg_one = MAGMA_D_NEG_ONE;
double *h_A, *h_R, *tau, *h_work, tmp[1];
double *d_lA[ MagmaMaxGPUs ];
magma_int_t M, N, n2, lda, ldda, n_local, ngpu;
magma_int_t info, min_mn, nb, lhwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1}, ISEED2[4];
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= (opts.check == 2); // check (-c2) implies lapack (-l)
magma_int_t status = 0;
double tol, eps = lapackf77_dlamch("E");
tol = opts.tolerance * eps;
printf("ngpu %d\n", (int) opts.ngpu );
if ( opts.check == 1 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R-Q'A||_1 / (M*||A||_1) ||I-Q'Q||_1 / M\n");
printf("================================================================================================\n");
} else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) ||R||_F /(M*||A||_F)\n");
printf("==========================================================================\n");
}
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
nb = magma_get_dgeqrf_nb( M );
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
// ngpu must be at least the number of blocks
ngpu = min( opts.ngpu, int((N+nb-1)/nb) );
if ( ngpu < opts.ngpu ) {
printf( " * too many GPUs for the matrix size, using %d GPUs\n", (int) ngpu );
}
// query for workspace size
lhwork = -1;
lapackf77_dgeqrf( &M, &N, h_A, &M, tau, tmp, &lhwork, &info );
lhwork = (magma_int_t) MAGMA_D_REAL( tmp[0] );
// Allocate host memory for the matrix
TESTING_MALLOC( tau, double, min_mn );
TESTING_MALLOC( h_A, double, n2 );
TESTING_HOSTALLOC( h_R, double, n2 );
TESTING_MALLOC( h_work, double, lhwork );
// Allocate device memory
for( int dev = 0; dev < ngpu; dev++ ) {
n_local = ((N/nb)/ngpu)*nb;
if (dev < (N/nb) % ngpu)
n_local += nb;
else if (dev == (N/nb) % ngpu)
n_local += N % nb;
magma_setdevice( dev );
TESTING_DEVALLOC( d_lA[dev], double, ldda*n_local );
}
/* Initialize the matrix */
for ( int j=0; j<4; j++ ) ISEED2[j] = ISEED[j]; // saving seeds
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
double *tau;
TESTING_MALLOC( tau, double, min_mn );
cpu_time = magma_wtime();
lapackf77_dgeqrf( &M, &N, h_A, &M, tau, h_work, &lhwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapack_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE( tau );
}
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix_1D_col_bcyclic( M, N, h_R, lda, d_lA, ldda, ngpu, nb );
gpu_time = magma_wtime();
magma_dgeqrf2_mgpu( ngpu, M, N, d_lA, ldda, tau, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgeqrf2 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_dgetmatrix_1D_col_bcyclic( M, N, d_lA, ldda, h_R, lda, ngpu, nb );
magma_queue_sync( NULL );
if ( opts.check == 1 ) {
/* =====================================================================
Check the result
=================================================================== */
magma_int_t lwork = n2+N;
double *h_W1, *h_W2, *h_W3;
double *h_RW, results[2];
TESTING_MALLOC( h_W1, double, n2 ); // Q
TESTING_MALLOC( h_W2, double, n2 ); // R
TESTING_MALLOC( h_W3, double, lwork ); // WORK
TESTING_MALLOC( h_RW, double, M ); // RWORK
lapackf77_dlarnv( &ione, ISEED2, &n2, h_A );
lapackf77_dqrt02( &M, &N, &min_mn, h_A, h_R, h_W1, h_W2, &lda, tau, h_W3, &lwork,
h_RW, results );
results[0] *= eps;
results[1] *= eps;
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %8.2e",
(int) M, (int) N, gpu_perf, gpu_time, results[0],results[1] );
printf("%s\n", (results[0] < tol ? "" : " failed"));
}
status |= ! (results[0] < tol);
TESTING_FREE( h_W1 );
TESTING_FREE( h_W2 );
TESTING_FREE( h_W3 );
TESTING_FREE( h_RW );
} else if ( opts.check == 2 ) {
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
error = lapackf77_dlange("f", &M, &N, h_A, &lda, work );
blasf77_daxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
error = lapackf77_dlange("f", &M, &N, h_R, &lda, work ) / (min_mn*error);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time, error );
printf("%s\n", (error < tol ? "" : " failed"));
status |= ! (error < tol);
}
else {
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) ---\n",
(int) M, (int) N, cpu_perf, cpu_time, gpu_perf, gpu_time );
} else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, gpu_time);
}
}
TESTING_FREE( tau );
TESTING_FREE( h_A );
TESTING_FREE( h_work );
TESTING_HOSTFREE( h_R );
for( int dev=0; dev < ngpu; dev++ ) {
magma_setdevice( dev );
TESTING_DEVFREE( d_lA[dev] );
}
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
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
=======
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 */
extern "C" magma_err_t
magma_cgeqr2x3_gpu(magma_int_t *m, magma_int_t *n,
magmaFloatComplex_ptr dA, size_t dA_offset, magma_int_t *ldda,
magmaFloatComplex_ptr dtau, size_t dtau_offset,
magmaFloatComplex_ptr dT, size_t dT_offset,
magmaFloatComplex_ptr ddA, size_t ddA_offset,
magmaFloat_ptr dwork, size_t dwork_offset,
magma_int_t *info, magma_queue_t queue)
{
/* -- clMAGMA (version 1.1.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date January 2014
Purpose
=======
CGEQR2 computes a QR factorization of a complex m by n matrix A:
A = Q * R.
This expert routine requires two more arguments than the standard
cgeqr2, namely, dT and ddA, explained below. The storage for A is
also not as in the LAPACK's cgeqr2 routine (see below).
The first is used to output the triangular
n x n factor T of the block reflector used in the factorization.
The second holds the diagonal nxn blocks of A, i.e., the diagonal
submatrices of R. This routine implements the left looking QR.
This version adds internal blocking.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the m by n matrix A.
On exit, the unitary matrix Q as a
product of elementary reflectors (see Further Details).
the elements on and above the diagonal of the array
contain the min(m,n) by n upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the unitary matrix Q as a
product of elementary reflectors (see Further Details).
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
dT (output) COMPLEX array, dimension N x N.
Stores the triangular N x N factor T of the block reflector
used in the factorization. The lower triangular part is 0.
ddA (output) COMPLEX array, dimension N x N.
Stores the elements of the upper N x N diagonal block of A.
LAPACK stores this array in A. There are 0s below the diagonal.
RWORK (workspace) DOUBLE_PRECISION array, dimension (3 N)
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
//#define da_ref(a_1,a_2) ( dA+(a_2)*(*ldda) + (a_1))
#define da_ref(a_1,a_2) dA, (dA_offset + ((a_2)*(*ldda) + (a_1)))
#define BLOCK_SIZE 32
//#define BLOCK_SIZE 16
static magma_int_t i, k;
//float *dnorm = dwork;
magmaFloat_ptr dnorm = dwork;
size_t dnorm_offset = dwork_offset;
//magmaFloatComplex *work = (magmaFloatComplex *)(dwork+2*(*n));
magmaFloatComplex_ptr work = (magmaFloatComplex_ptr)dwork;
size_t work_offset = dwork_offset + 2*(*n);
*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;
}
/* Compute the norms of the trailing columns */
k = min(*m,*n);
magmablas_scnrm2(*m, k, da_ref(0,0), *ldda, dnorm, dnorm_offset, queue);
for (int b=0; b < k; b += BLOCK_SIZE) {
for (i = b; i < min(k, b+BLOCK_SIZE); ++i) {
/* Apply H' to A(:,i) from the left */
if ( i-b > 0){
magma_queue_sync(queue);
magma_clarfbx_gpu(*m-b, i-b, da_ref(b, b), *ldda,
dT, (dT_offset+b+b*k), k, da_ref(b, i), work, work_offset, queue);
}
/* Adjust the dnorm[i] to hold the norm of A(i:m,i) */
if ( i > 0 ){
magma_queue_sync(queue);
magmablas_scnrm2_adjust(i, dnorm, dnorm_offset+i, da_ref(0, i), queue);
}
/* Generate elementary reflector H(i) to annihilate A(i+1:m,i)
1. 1 is not yet put on the diagonal of A
2. Elements above the diagonal are copied in ddA and
the ones in A are set to zero
3. update T */
magma_clarfgtx_gpu(*m-i, da_ref(i, i), da_ref(min(i+1,*m), i), dtau, dtau_offset+i,
dnorm, dnorm_offset+i, ddA, ddA_offset + i + i*(*n), i,
da_ref(i,0), *ldda, dT, dT_offset, k, work, work_offset, queue);
}
/* Apply the transformations to the trailing matrix. */
magma_clarfb2_gpu(
*m-b, k-i, BLOCK_SIZE,
da_ref(b, b), *ldda, dT, dT_offset+b+b*k, k,
da_ref(b, i), *ldda, work, work_offset, k-i, queue);
}
magma_queue_sync(queue);
return *info;
} /* magma_cgeqr2 */
extern "C" magma_int_t
magma_sgeqrf2_mgpu( magma_int_t num_gpus, magma_int_t m, magma_int_t n,
float **dlA, magma_int_t ldda,
float *tau,
magma_int_t *info )
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SGEQRF2_MGPU computes a QR factorization of a real M-by-N matrix A:
A = Q * R. This is a GPU interface of the routine.
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.
dA (input/output) REAL array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix dA.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
TAU (output) REAL array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define dlA(gpu,a_1,a_2) ( dlA[gpu]+(a_2)*(ldda) + (a_1))
#define work_ref(a_1) ( work + (a_1))
#define hwork ( work + (nb)*(m))
#define hwrk_ref(a_1) ( local_work + (a_1))
#define lhwrk ( local_work + (nb)*(m))
float *dwork[4], *panel[4], *local_work;
magma_int_t i, j, k, ldwork, lddwork, old_i, old_ib, rows;
magma_int_t nbmin, nx, ib, nb;
magma_int_t lhwork, lwork;
magma_device_t cdevice;
magma_getdevice(&cdevice);
int panel_gpunum, i_local, n_local[4], la_gpu, displacement;
*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;
}
k = min(m,n);
if (k == 0)
return *info;
nb = magma_get_sgeqrf_nb(m);
displacement = n * nb;
lwork = (m+n+64) * nb;
lhwork = lwork - (m)*nb;
for(i=0; i<num_gpus; i++){
#ifdef MultiGPUs
magma_setdevice(i);
#endif
if (MAGMA_SUCCESS != magma_smalloc( &(dwork[i]), (n + ldda)*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
/* Set the number of local n for each GPU */
for(i=0; i<num_gpus; 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;
}
if (MAGMA_SUCCESS != magma_smalloc_pinned( &local_work, lwork )) {
*info = -9;
for(i=0; i<num_gpus; i++){
#ifdef MultiGPUs
magma_setdevice(i);
#endif
magma_free( dwork[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
cudaStream_t streaml[4][2];
for(i=0; i<num_gpus; i++){
#ifdef MultiGPUs
magma_setdevice(i);
#endif
magma_queue_create( &streaml[i][0] );
magma_queue_create( &streaml[i][1] );
}
nbmin = 2;
nx = nb;
ldwork = m;
lddwork= n;
if (nb >= nbmin && nb < k && nx < k) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nx; i += nb)
{
/* Set the GPU number that holds the current panel */
panel_gpunum = (i/nb)%num_gpus;
/* Set the local index where the current panel is */
i_local = i/(nb*num_gpus)*nb;
ib = min(k-i, nb);
rows = m -i;
/* Send current panel to the CPU */
#ifdef MultiGPUs
magma_setdevice(panel_gpunum);
#endif
magma_sgetmatrix_async( rows, ib,
dlA(panel_gpunum, i, i_local), ldda,
hwrk_ref(i), ldwork, streaml[panel_gpunum][1] );
if (i>0){
/* Apply H' to A(i:m,i+2*ib:n) from the left; this is the look-ahead
application to the trailing matrix */
la_gpu = panel_gpunum;
/* only the GPU that has next panel is done look-ahead */
#ifdef MultiGPUs
magma_setdevice(la_gpu);
#endif
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
m-old_i, n_local[la_gpu]-i_local-old_ib, old_ib,
panel[la_gpu], ldda, dwork[la_gpu], lddwork,
dlA(la_gpu, old_i, i_local+old_ib), ldda,
dwork[la_gpu]+old_ib, lddwork);
la_gpu = ((i-nb)/nb)%num_gpus;
#ifdef MultiGPUs
magma_setdevice(la_gpu);
#endif
magma_ssetmatrix_async( old_ib, old_ib,
hwrk_ref(old_i), ldwork,
panel[la_gpu], ldda, streaml[la_gpu][0] );
}
#ifdef MultiGPUs
magma_setdevice(panel_gpunum);
#endif
magma_queue_sync( streaml[panel_gpunum][1] );
lapackf77_sgeqrf(&rows, &ib, hwrk_ref(i), &ldwork, tau+i, lhwrk, &lhwork, info);
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
hwrk_ref(i), &ldwork, tau+i, lhwrk, &ib);
spanel_to_q( MagmaUpper, ib, hwrk_ref(i), ldwork, lhwrk+ib*ib );
// Send the current panel back to the GPUs
// Has to be done with asynchronous copies
for(j=0; j<num_gpus; j++)
{
#ifdef MultiGPUs
magma_setdevice(j);
#endif
if (j == panel_gpunum)
panel[j] = dlA(j, i, i_local);
else
panel[j] = dwork[j]+displacement;
magma_ssetmatrix_async( rows, ib,
hwrk_ref(i), ldwork,
panel[j], ldda, streaml[j][0] );
}
for(j=0; j<num_gpus; j++)
{
#ifdef MultiGPUs
magma_setdevice(j);
#endif
magma_queue_sync( streaml[j][0] );
}
/* Restore the panel */
sq_to_panel( MagmaUpper, ib, hwrk_ref(i), ldwork, lhwrk+ib*ib );
if (i + ib < n)
{
/* Send the T matrix to the GPU.
Has to be done with asynchronous copies */
for(j=0; j<num_gpus; j++)
{
#ifdef MultiGPUs
magma_setdevice(j);
#endif
magma_ssetmatrix_async( ib, ib,
lhwrk, ib,
dwork[j], lddwork, streaml[j][0] );
}
if (i+nb < k-nx)
{
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left;
This is update for the next panel; part of the look-ahead */
la_gpu = (panel_gpunum+1)%num_gpus;
int i_loc = (i+nb)/(nb*num_gpus)*nb;
for(j=0; j<num_gpus; j++){
#ifdef MultiGPUs
magma_setdevice(j);
#endif
//magma_queue_sync( streaml[j][0] );
if (j==la_gpu)
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
panel[j], ldda, dwork[j], lddwork,
dlA(j, i, i_loc), ldda, dwork[j]+ib, lddwork);
else if (j<=panel_gpunum)
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
rows, n_local[j]-i_local-ib, ib,
panel[j], ldda, dwork[j], lddwork,
dlA(j, i, i_local+ib), ldda, dwork[j]+ib, lddwork);
else
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
rows, n_local[j]-i_local, ib,
panel[j], ldda, dwork[j], lddwork,
dlA(j, i, i_local), ldda, dwork[j]+ib, lddwork);
}
}
else {
/* do the entire update as we exit and there would be no lookahead */
la_gpu = (panel_gpunum+1)%num_gpus;
int i_loc = (i+nb)/(nb*num_gpus)*nb;
#ifdef MultiGPUs
magma_setdevice(la_gpu);
#endif
magma_slarfb_gpu( MagmaLeft, MagmaTrans, MagmaForward, MagmaColumnwise,
rows, n_local[la_gpu]-i_loc, ib,
panel[la_gpu], ldda, dwork[la_gpu], lddwork,
dlA(la_gpu, i, i_loc), ldda, dwork[la_gpu]+ib, lddwork);
#ifdef MultiGPUs
magma_setdevice(panel_gpunum);
#endif
magma_ssetmatrix( ib, ib,
hwrk_ref(i), ldwork,
dlA(panel_gpunum, i, i_local), ldda );
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
for(j=0; j<num_gpus; j++){
#ifdef MultiGPUs
magma_setdevice(j);
#endif
magma_free( dwork[j] );
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
lhwork = lwork - rows*ib;
panel_gpunum = (panel_gpunum+1)%num_gpus;
int i_loc = (i)/(nb*num_gpus)*nb;
#ifdef MultiGPUs
magma_setdevice(panel_gpunum);
#endif
magma_sgetmatrix( rows, ib,
dlA(panel_gpunum, i, i_loc), ldda,
lhwrk, rows );
lhwork = lwork - rows*ib;
lapackf77_sgeqrf(&rows, &ib, lhwrk, &rows, tau+i, lhwrk+ib*rows, &lhwork, info);
magma_ssetmatrix( rows, ib,
lhwrk, rows,
dlA(panel_gpunum, i, i_loc), ldda );
}
for(i=0; i<num_gpus; i++){
#ifdef MultiGPUs
magma_setdevice(i);
#endif
magma_queue_destroy( streaml[i][0] );
magma_queue_destroy( streaml[i][1] );
}
magma_setdevice(cdevice);
magma_free_pinned( local_work );
return *info;
} /* magma_sgeqrf2_mgpu */
extern "C" magma_int_t
magma_sgeqrf_msub(
magma_int_t num_subs, magma_int_t num_gpus,
magma_int_t m, magma_int_t n,
magmaFloat_ptr *dlA, magma_int_t ldda,
float *tau,
magma_queue_t *queues,
magma_int_t *info)
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
SGEQRF2_MGPU computes a QR factorization of a real M-by-N matrix A:
A = Q * R. This is a GPU interface of the routine.
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.
dA (input/output) REAL array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix dA.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
TAU (output) REAL array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
Further Details
===============
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
===================================================================== */
#define dlA(gpu,a_1,a_2) dlA[gpu], ((a_2)*(ldda) + (a_1))
#define dlA_offset(a_1, a_2) ((a_2)*(ldda) + (a_1))
#define work_ref(a_1) ( work + (a_1))
#define hwork ( work + (nb)*(m))
#define hwrk(a_1) ( local_work + (a_1))
#define lhwrk ( local_work + (nb)*(m))
magmaFloat_ptr dwork[MagmaMaxGPUs], panel[MagmaMaxGPUs];
size_t panel_offset[MagmaMaxGPUs];
float *local_work = NULL;
magma_int_t i, j, k, ldwork, lddwork, old_i, old_ib, rows;
magma_int_t nbmin, nx, ib, nb;
magma_int_t lhwork, lwork;
int panel_id = -1, i_local, n_local[MagmaMaxGPUs * MagmaMaxSubs], la_id, displacement,
tot_subs = num_gpus * num_subs;
*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;
}
k = min(m,n);
if (k == 0)
return *info;
nb = magma_get_sgeqrf_nb(m);
displacement = n * nb;
lwork = (m+n+64) * nb;
lhwork = lwork - (m)*nb;
for (i=0; i<num_gpus; i++) {
if (MAGMA_SUCCESS != magma_smalloc( &(dwork[i]), (n + ldda)*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
/* Set the number of local n for each GPU */
for (i=0; i<tot_subs; i++) {
n_local[i] = ((n/nb)/tot_subs)*nb;
if (i < (n/nb)%tot_subs)
n_local[i] += nb;
else if (i == (n/nb)%tot_subs)
n_local[i] += n%nb;
}
#ifdef USE_PINNED_CLMEMORY
cl_mem buffer = clCreateBuffer(gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(float)*lwork, NULL, NULL);
for (j=0; j<num_gpus; j++) {
local_work = (float*)clEnqueueMapBuffer(queues[2*j], buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0,
sizeof(float)*lwork, 0, NULL, NULL, NULL);
}
#else
if (MAGMA_SUCCESS != magma_smalloc_cpu( (&local_work), lwork )) {
*info = -9;
for (i=0; i<num_gpus; i++) {
magma_free( dwork[i] );
}
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
nbmin = 2;
nx = nb;
ldwork = m;
lddwork= n;
if (nb >= nbmin && nb < k && nx < k) {
/* Use blocked code initially */
old_i = 0; old_ib = nb;
for (i = 0; i < k-nx; i += nb) {
/* Set the GPU number that holds the current panel */
panel_id = (i/nb)%tot_subs;
/* Set the local index where the current panel is */
i_local = i/(nb*tot_subs)*nb;
ib = min(k-i, nb);
rows = m -i;
/* Send current panel to the CPU */
magma_queue_sync(queues[2*(panel_id%num_gpus)]);
magma_sgetmatrix_async( rows, ib,
dlA(panel_id, i, i_local), ldda,
hwrk(i), ldwork,
queues[2*(panel_id%num_gpus)+1], NULL );
if (i > 0) {
/* Apply H' to A(i:m,i+2*ib:n) from the left; this is the look-ahead
application to the trailing matrix */
la_id = panel_id;
/* only the GPU that has next panel is done look-ahead */
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n_local[la_id]-i_local-old_ib, old_ib,
panel[la_id%num_gpus], panel_offset[la_id%num_gpus], ldda,
dwork[la_id%num_gpus], 0, lddwork,
dlA(la_id, old_i, i_local+old_ib), ldda,
dwork[la_id%num_gpus], old_ib, lddwork,
queues[2*(la_id%num_gpus)]);
la_id = ((i-nb)/nb)%tot_subs;
magma_ssetmatrix_async( old_ib, old_ib,
hwrk(old_i), ldwork,
panel[la_id%num_gpus], panel_offset[la_id%num_gpus], ldda,
queues[2*(la_id%num_gpus)], NULL );
}
magma_queue_sync( queues[2*(panel_id%num_gpus)+1] );
lapackf77_sgeqrf(&rows, &ib, hwrk(i), &ldwork, tau+i, lhwrk, &lhwork, info);
// Form the triangular factor of the block reflector
// H = H(i) H(i+1) . . . H(i+ib-1)
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib,
hwrk(i), &ldwork, tau+i, lhwrk, &ib);
spanel_to_q( MagmaUpper, ib, hwrk(i), ldwork, lhwrk+ib*ib );
// Send the current panel back to the GPUs
// Has to be done with asynchronous copies
for (j=0; j<num_gpus; j++) {
if (j == panel_id%num_gpus){
panel[j] = dlA(panel_id, i, i_local);
panel_offset[j] = dlA_offset(i, i_local);
} else {
panel[j] = dwork[j];
panel_offset[j] = displacement;
}
magma_queue_sync( queues[2*j] );
magma_ssetmatrix_async( rows, ib,
hwrk(i), ldwork,
panel[j], panel_offset[j], ldda,
queues[2*j+1], NULL );
/* Send the T matrix to the GPU.
Has to be done with asynchronous copies */
magma_ssetmatrix_async( ib, ib, lhwrk, ib,
dwork[j], 0, lddwork,
queues[2*j+1], NULL );
}
for(j=0; j<num_gpus; j++) {
magma_queue_sync( queues[2*j+1] );
}
if (i + ib < n) {
if (i+nb < k-nx) {
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left;
This is update for the next panel; part of the look-ahead */
la_id = (panel_id+1)%tot_subs;
int i_loc = (i+nb)/(nb*tot_subs)*nb;
for (j=0; j<tot_subs; j++) {
if (j == la_id)
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
panel[j%num_gpus], panel_offset[j%num_gpus], ldda,
dwork[j%num_gpus], 0, lddwork,
dlA(j, i, i_loc), ldda,
dwork[j%num_gpus], ib, lddwork,
queues[2*(j%num_gpus)]);
else if (j <= panel_id)
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[j]-i_local-ib, ib,
panel[j%num_gpus], panel_offset[j%num_gpus], ldda,
dwork[j%num_gpus], 0, lddwork,
dlA(j, i, i_local+ib), ldda,
dwork[j%num_gpus], ib, lddwork,
queues[2*(j%num_gpus)]);
else
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[j]-i_local, ib,
panel[j%num_gpus], panel_offset[j%num_gpus], ldda,
dwork[j%num_gpus], 0, lddwork,
dlA(j, i, i_local), ldda,
dwork[j%num_gpus], ib, lddwork,
queues[2*(j%num_gpus)]);
}
/* Restore the panel */
sq_to_panel( MagmaUpper, ib, hwrk(i), ldwork, lhwrk+ib*ib );
} else {
/* do the entire update as we exit and there would be no lookahead */
la_id = (panel_id+1)%tot_subs;
int i_loc = (i+nb)/(nb*tot_subs)*nb;
magma_slarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n_local[la_id]-i_loc, ib,
panel[la_id%num_gpus], panel_offset[la_id%num_gpus], ldda,
dwork[la_id%num_gpus], 0, lddwork,
dlA(la_id, i, i_loc), ldda,
dwork[la_id%num_gpus], ib, lddwork,
queues[2*(la_id%num_gpus)]);
/* Restore the panel */
sq_to_panel( MagmaUpper, ib, hwrk(i), ldwork, lhwrk+ib*ib );
magma_ssetmatrix( ib, ib,
hwrk(i), ldwork,
dlA(panel_id, i, i_local), ldda,
queues[2*(panel_id%num_gpus)]);
}
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
for (j=0; j<num_gpus; j++) {
magma_free( dwork[j] );
}
/* Use unblocked code to factor the last or only block. */
if (i < k) {
ib = n-i;
rows = m-i;
lhwork = lwork - rows*ib;
panel_id = (panel_id+1)%tot_subs;
int i_loc = (i)/(nb*tot_subs)*nb;
magma_sgetmatrix( rows, ib,
dlA(panel_id, i, i_loc), ldda,
lhwrk, rows,
queues[2*(panel_id%num_gpus)]);
lhwork = lwork - rows*ib;
lapackf77_sgeqrf(&rows, &ib, lhwrk, &rows, tau+i, lhwrk+ib*rows, &lhwork, info);
magma_ssetmatrix( rows, ib,
lhwrk, rows,
dlA(panel_id, i, i_loc), ldda,
queues[2*(panel_id%num_gpus)]);
}
#ifdef USE_PINNED_CLMEMORY
#else
magma_free_cpu( local_work );
#endif
return *info;
} /* magma_sgeqrf_msub */