void magmaf_dlarfb_gpu_gemm(
magma_side_t *side, magma_trans_t *trans, magma_direct_t *direct, magma_storev_t *storev, magma_int_t *m, magma_int_t *n, magma_int_t *k,
const double *dv, magma_int_t *ldv,
const double *dt, magma_int_t *ldt,
double *dc, magma_int_t *ldc,
double *dwork, magma_int_t *ldwork,
double *dworkvt, magma_int_t *ldworkvt )
{
magma_dlarfb_gpu_gemm(
*side, *trans, *direct, *storev, *m, *n, *k,
dv, *ldv,
dt, *ldt,
dc, *ldc,
dwork, *ldwork,
dworkvt, *ldworkvt );
}
/***************************************************************************//**
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;
}
extern "C" magma_int_t
magma_dbulge_applyQ_v2_m(
magma_int_t ngpu,
magma_side_t side,
magma_int_t NE, magma_int_t N,
magma_int_t NB, magma_int_t Vblksiz,
double *E, magma_int_t lde,
double *V, magma_int_t ldv,
double *T, magma_int_t ldt,
magma_int_t *info)
{
//%===========================
//% local variables
//%===========================
magma_int_t Vm, Vn, mt, nt;
magma_int_t myrow, mycol, blkj, blki;
magma_int_t blkid,vpos,tpos;
magma_int_t firstrow, nbcolinvolvd;
magma_int_t versionL = 113;
magma_int_t versionR = 92;
magma_int_t Vchunksiz = 10;
*info=0;
/* Quick return */
if ( NE == 0 ) {
return *info;
}
if ( N == 0 ) {
return *info;
}
if ( NB == 0 ) {
return *info;
}
/* ==========================================
* some infos for developer
* Initialisation and checking nb of cores
* ==========================================*/
/* we have 2 algo for left (113 114) and 2 algo for right (91 92)
* which correspond to versionL versionR.
* They are very similar (detail explained in tech report and matlab code)
* however version 114 and 92 improve locality.
* while version 113 is used in case WNATZ=1 (construct Q2) which allow
* the construction to be done in an optimized way taking into
* consideration that the matrix is Identity so making less flops.
*
*/
// Initialize streaming and events
//magma_device_sync();
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_int_t nevents =2, nstream=2;
magma_queue_t queues[MagmaMaxGPUs][20];
magma_event_t myevent[MagmaMaxGPUs][20];
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( dev, &queues[dev][i] );
}
for( magma_int_t i = 0; i < nevents; ++i ) {
cudaEventCreateWithFlags(&myevent[dev][i],cudaEventDisableTiming);
//magma_event_create(&myevent[dev][i]);
}
}
// Azzam 21/11/2012
// NOTE THAT dwork was of size 2*NE*Vblksiz+...
// but I am thinking why not modifing it to NE*Vblksiz+...
// BUT NO because the 2* is used because of making 2 queues working and so
// they might be using dwork in parallel
double *dE[MagmaMaxGPUs];
double *dwork[MagmaMaxGPUs], *dwork0[MagmaMaxGPUs], *dwork1[MagmaMaxGPUs];
//double *dwvt[MagmaMaxGPUs];
double *dwvt0[MagmaMaxGPUs], *dwvt1[MagmaMaxGPUs];
double *dT0[MagmaMaxGPUs], *dV0[MagmaMaxGPUs], *dT1[MagmaMaxGPUs], *dV1[MagmaMaxGPUs];
magma_int_t dev;
magma_int_t ldde = N;
magma_int_t lddv = ldv;
magma_int_t lddt = ldt;
magma_int_t ne_loc = magma_ceildiv(NE, ngpu);
if (ne_loc < 256)
ne_loc=256;
magma_int_t dwVTsiz = lddv*Vblksiz; // lddv*lddv + lddv*NE; // lddv*Vblksiz;
magma_int_t dworksiz = ne_loc*Vblksiz; // lddv*Vblksiz; // NE*Vblksiz;
ngpu = min(ngpu, magma_ceildiv(NE,ne_loc)); // Don't use GPU that will not have data.
// copy dE to GPUs
for (dev=0; dev < ngpu; ++dev) {
magma_setdevice( dev );
if (MAGMA_SUCCESS != magma_dmalloc( &dE[dev], ldde * ne_loc)) {
printf ("!!!! magma_dbulge_applyQ magma_alloc failed for: dE\n" );
exit(-1);
}
if (MAGMA_SUCCESS != magma_dmalloc( &dwork[dev], 2*dworksiz + 2*dwVTsiz + 2*Vchunksiz* (Vblksiz* (lddv+lddt)) )) {
printf ("!!!! magma_dbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
dwork0[dev] = dwork[dev]; // size = dworksiz;
dwork1[dev] = dwork0[dev] + dworksiz; // size = dworksiz;
dwvt0[dev] = dwork[dev] + 2*dworksiz; // size = dwVTsiz;
dwvt1[dev] = dwvt0[dev] + dwVTsiz; // size = dwVTsiz;
dV0[dev] = dwork[dev] + 2*dworksiz + 2*dwVTsiz;
dT0[dev] = dV0[dev] + Vchunksiz*Vblksiz*lddv;
dV1[dev] = dT0[dev] + Vchunksiz*Vblksiz*lddt;
dT1[dev] = dV1[dev] + Vchunksiz*Vblksiz*lddv;
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_dsetmatrix_async( N, ie_loc, E+lde*ne_loc*dev, lde, dE(dev, 0, 0), ldde, queues[dev][1] );
}
// make overlapped copy
magma_int_t ncpy = 0;
magma_int_t copyed=0, copyst=0;
magma_int_t blkcnt,nothing, mysiz, flip, vld,tld, locpos;
findVTsiz(N, NB, Vblksiz, &blkcnt, ¬hing);
flip = 0;
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
#ifdef ENABLE_DEBUG
printf(" APPLY Q_v22_m GPU with NGPU %d N %d, NE %d, NB %d, Vblksiz %d, versionL %d versionR %d SIDE %c \n",
ngpu, N, NE, NB, Vblksiz, versionL, versionR, side);
#endif
/*
* MagmamaLeft
*/
if (side == MagmaLeft) {
/*
* Version 113:
* loop over the block_col (nt) and for each find the
* number of tiles (mt) in this block_col. then loop over mt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
if ( versionL == 113 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=nt-1; blkj >= 0; blkj--) {
/* the index of the first row on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=mt; blki > 0; blki--) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d vpos %d \n",blkj,blki,Vm, Vn,mycol,vpos);
magma_bulge_findpos113(N, NB, Vblksiz, mycol, myrow, &blkid);
// COPY Vchunksiz Vs and Vchunksiz Ts to GPU and store it in dV0/dV1 and dT0/dT1
if (ncpy == 0) {
// flip = 1 for this.
copyst = 0; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
if (mysiz > 0) {
ncpy = 1;
flip = 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_dsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, queues[dev][1]);
magma_dsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, queues[dev][1]);
}
//printf("doing the first copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
}
}
if (blkid == copyst) {
flip = ncpy % 2;
copyst = copyed; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
//printf(" get to copy blkid %d blkid+(2*Vchunksiz) %d copyst %d copyed %d\n",blkid,blkid+(Vchunksiz),copyst,copyed);
if (mysiz > 0) {
ncpy = ncpy + 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
if (flip == 0) { // now I am working on dV0 so copy the next and put it on dV1
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_dsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, queues[dev][1]);
magma_dsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, queues[dev][1]);
}
} else { // now I am working on dV1 so copy the next and put it on dV0
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV0 dT0\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_dsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV0[dev], vld, queues[dev][0]);
magma_dsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT0[dev], tld, queues[dev][0]);
}
}
}
}
if ((Vm > 0) && (Vn > 0)) {
locpos = blkid%Vchunksiz;
magma_int_t lcvpos = locpos*Vblksiz*lddv;
magma_int_t lctpos = locpos*Vblksiz*lddt;
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d locvpos %5d loctpos %5d blkid %2d using data in dV%1d dT%1d \n",blkj,blki,Vm, Vn,mycol,lcvpos,lctpos, blkid,flip,flip);
if (flip == 0) {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magma_queue_wait_event( queues[dev][0], myevent[dev][1] );
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib, queues[dev][0] );
magma_dlarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib, dwvt0[dev], Vm, queues[dev][0] );
}
magma_event_record( myevent[dev][0], queues[dev][0] );
}
} else {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magma_queue_wait_event( queues[dev][1], myevent[dev][0] );
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib, queues[dev][1] );
magma_dlarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib, dwvt1[dev], Vm, queues[dev][1] );
}
magma_event_record( myevent[dev][1], queues[dev][1] );
}
}
} // end for (Vm &Vn) > 0
} // end for blki
} // end for blkj
} // end if version=113
/*
* Version 114:
* loop over the block_row (mt) and for each find diagonally the
* number of tiles (nt) in this block_row. then loop over nt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
else {
printf("versionL 114 not implemented in dbulge_applyQ_v2_m\n");
exit(-1);
mt = magma_ceildiv((N-1),NB);
for (blki = mt; blki > 0; blki--) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = nt-1; blkj >= 0; blkj--) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_dsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, queues[dev][0]);
magma_dsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, queues[dev][0]);
//printf("voici blki %d rownbm %d mycol %d coled %d blkid %d vpos %d tpos %d\n", blki, rownbm, mycol, coled, blkid, vpos, tpos);
for (magma_int_t i=0; i < NE; i += sz_bl) {
ib = min(sz_bl, NE-i);
magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0, lddv, dT0, lddt, dE(myrow,i), ldde, dwork, NE, queues[dev][0] );
}
*/
} // end for (Vm &Vn) > 0
} // end for blkj
} // end for blki
} // end version 114
} // end LEFT
/*
* MagmaRight
*/
else {
printf("Side 'R' not implemented in dbulge_applyQ_v2_m\n");
exit(-1);
/*
* Version 91:
*/
if ( versionR == 91 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=0; blkj < nt; blkj++) {
/* the index of the first myrow on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=1; blki <= mt; blki++) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( (blkj == nt-1) && (blki == mt) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
mycol = blkj*Vblksiz;
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_dsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, queues[dev][0]);
magma_dsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, queues[dev][0]);
magma_dlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE, queues[dev][0] );
*/
} // end for (Vm &Vn) > 0
} // end for blki
} // end fo blkj
} // end of version 91
/*
* Version 92:
*/
else {
mt = magma_ceildiv((N-1),NB);
for (blki = 1; blki <= mt; blki++) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = 0; blkj < nt; blkj++) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_dsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, queues[dev][0]);
magma_dsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, queues[dev][0]);
magma_dlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE, queues[dev][0] );
*/
} // end for (Vm &Vn) > 0
} //end for blkj
} // end for blki
} //end of version 92
} // end RIGHT
// copy back the dE form each GPU
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_wait_event( queues[dev][0], myevent[dev][1] );
magma_queue_wait_event( queues[dev][0], myevent[dev][0] );
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_dgetmatrix_async( N, ie_loc, dE(dev, 0, 0), ldde, E+lde*ne_loc*dev, lde, queues[dev][0] );
magma_event_record( myevent[dev][0], queues[dev][0] );
}
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_wait_event( queues[dev][0], myevent[dev][0] );
magma_queue_sync(queues[dev][0]);
magma_queue_sync(queues[dev][1]);
magma_free(dwork[dev]);
magma_free(dE[dev]);
for( magma_int_t i = 0; i < nevents; ++i ) {
magma_event_destroy( myevent[dev][i] );
}
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( queues[dev][i] );
}
}
magma_setdevice( orig_dev );
return *info;
}
