extern "C" magma_int_t
magma_sgeqrf_expert_batched(
magma_int_t m, magma_int_t n,
float **dA_array, magma_int_t ldda,
float **dR_array, magma_int_t lddr,
float **dT_array, magma_int_t lddt,
float **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) // A(i, j) means at i row, j column
/* Local Parameter */
magma_int_t nb = magma_get_sgeqrf_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;
float **dW0_displ = NULL;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dW4_displ = NULL;
float **dW5_displ = NULL;
float **dR_displ = NULL;
float **dT_displ = NULL;
float *dwork = NULL;
float **cpuAarray = NULL;
float **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_smalloc(&dwork, (2 * nb * n) * batchCount);
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
magma_malloc_cpu((void**) &cpuTarray, batchCount*sizeof(float*));
/* 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_sdisplace_pointers(dR_displ, dR_array, lddr, 0, 0, batchCount, queue);
magma_sdisplace_pointers(dT_displ, dT_array, lddt, 0, 0, batchCount, 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_slaset_batched( MagmaFull, lddr, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_S_ZERO, MAGMA_S_ZERO, dR_displ, lddr, batchCount, queue );
magmablas_slaset_batched( MagmaFull, lddt, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_S_ZERO, MAGMA_S_ZERO, dT_displ, lddt, batchCount, queue );
/*
if ( provide_RT > 0 )
{
magmablas_slaset_q( MagmaFull, lddr, n*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dR, lddr, queue );
magmablas_slaset_q( MagmaFull, lddt, n*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dT, lddt, queue );
}
else
{
magmablas_slaset_q( MagmaFull, lddr, nb*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dR, lddr, queue );
magmablas_slaset_q( MagmaFull, lddt, nb*batchCount, MAGMA_S_ZERO, MAGMA_S_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(float*), dA_array, 1, cpuAarray, 1, queue);
magma_getvector( batchCount, sizeof(float*), dT_array, 1, cpuTarray, 1, queue);
for (i=0; i < min_mn; i += nb)
{
ib = min(nb, min_mn-i);
//===============================================
// panel factorization
//===============================================
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
if ( provide_RT > 0 )
{
offset_RT = i;
magma_sdisplace_pointers(dR_displ, dR_array, lddr, (provide_RT == 1 ? offset_RT:0), offset_RT, batchCount, queue);
magma_sdisplace_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_sgeqrf_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_sset_pointer( dW4_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
offset = ldw*n*batchCount;
magma_sset_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_slaset_batched( MagmaUpper, ib, ib, MAGMA_S_ZERO, MAGMA_S_ONE, dW0_displ, ldda, batchCount, queue );
//magma_sdisplace_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_slarft_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_srecommend_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_slarfb_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_sdisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_slarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
(const float**)dW0_displ, ldda,
(const float**)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_slacpy_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_sbulge_applyQ_v2(
magma_side_t side,
magma_int_t NE, magma_int_t N,
magma_int_t NB, magma_int_t Vblksiz,
magmaFloat_ptr dE, magma_int_t ldde,
float *V, magma_int_t ldv,
float *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_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_t myevent[2];
cudaEventCreateWithFlags(&myevent[0],cudaEventDisableTiming);
cudaEventCreateWithFlags(&myevent[1],cudaEventDisableTiming);
// Azzam 21/11/2012
// NOTE THAT dwork was of size 2*NE*Vblksiz+...
// but I am thinking why not modifing it to NE*Vblksiz+...
// BUT NO because the 2* is used because of making 2 streams working and so
// they might be using dwork in parallel
float *dwork, *dwork0, *dwork1, *dwvt0, *dwvt1;
float *dT0, *dV0, *dT1, *dV1;
magma_int_t lddv = ldv;
magma_int_t lddt = ldt;
magma_int_t lddw = 0;
magma_int_t lddwork = ((NE+31)/32)*32;
magma_int_t dwVTsiz = lddv*Vblksiz; // lddv*lddv + lddv*lddwork; (v2) // lddv*Vblksiz; (v1,v3)
magma_int_t dworksiz = lddwork*Vblksiz; // lddv*Vblksiz; (v2) // NE*Vblksiz=lddwork*Vblksiz; (v1,v3)
if (MAGMA_SUCCESS != magma_smalloc( &dwork, 2*dworksiz + 2*dwVTsiz + 2*Vchunksiz* (Vblksiz* (lddv+lddt)) )) {
printf ("!!!! magma_sbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
dwork0 = dwork; // size = dworksiz;
dwork1 = dwork0 + dworksiz; // size = dworksiz;
dwvt0 = dwork + 2*dworksiz; // size = dwVTsiz;
dwvt1 = dwvt0 + dwVTsiz; // size = dwVTsiz;
dV0 = dwork + 2*dworksiz + 2*dwVTsiz;
dT0 = dV0 + Vchunksiz*Vblksiz*lddv;
dV1 = dT0 + Vchunksiz*Vblksiz*lddt;
dT1 = dV1 + Vchunksiz*Vblksiz*lddv;
// 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;
// performance loss if the reflector are applied to a big number of eigenvectors (~10000)
// => apply the reflectors to blocks of eigenvectors.
//magma_int_t nr_bl = magma_ceildiv(NE,10000); //nr of blocks
magma_int_t sz_bl = NE; //magma_ceildiv(NE,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
/* 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 GPU with N %d, NE %d, NB %d, Vblksiz %d, versionL %d versionR %d SIDE %c \n",
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;
magmablasSetKernelStream(stream[1]);
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1, vld, stream[1]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1, tld, stream[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);
magmablasSetKernelStream(stream[1]);
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1, vld, stream[1]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1, tld, stream[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);
magmablasSetKernelStream(stream[0]);
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV0, vld, stream[0]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT0, tld, stream[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) {
magmablasSetKernelStream(stream[0]);
magma_queue_wait_event( stream[0], myevent[1] );
for (magma_int_t i=0; i < NE; i += sz_bl) {
ib = min(sz_bl, NE-i);
lddw = min(lddwork,sz_bl);
//magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0+lcvpos, lddv, dT0+lctpos, lddt, dE(myrow,i), ldde, dwork0, lddw);
magma_slarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0+lcvpos, lddv, dT0+lctpos, lddt, dE(myrow,i), ldde, dwork0, lddw, dwvt0, lddv);
}
magma_event_record( myevent[0], stream[0] );
} else {
magmablasSetKernelStream(stream[1]);
magma_queue_wait_event( stream[1], myevent[0] );
for (magma_int_t i=0; i < NE; i += sz_bl) {
ib = min(sz_bl, NE-i);
lddw = min(lddwork,sz_bl);
//magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1+lcvpos, lddv, dT1+lctpos, lddt, dE(myrow,i), ldde, dwork1, lddw);
magma_slarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1+lcvpos, lddv, dT1+lctpos, lddt, dE(myrow,i), ldde, dwork1, lddw, dwvt1, lddv);
}
magma_event_record( myevent[1], stream[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 {
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_ssetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_ssetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
//printf("voici blki %d rownbm %d mycol %d coled %d blkid %d vpos %d tpos %d\n", blki, rownbm, mycol, coled, blkid, vpos, tpos);
for (magma_int_t i=0; i < NE; i += sz_bl) {
ib = min(sz_bl, NE-i);
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0, lddv, dT0, lddt, dE(myrow,i), ldde, dwork, NE);
}
} // end for (Vm &Vn) > 0
} // end for blkj
} // end for blki
} // end version 114
} // end LEFT
/*
* MagmaRight
*/
else {
/*
* 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_ssetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_ssetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_slarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
} // end for (Vm &Vn) > 0
} // end for blki
} // end fo blkj
} // end of version 91
/*
* Version 92:
*/
else {
mt = magma_ceildiv((N-1),NB);
for (blki = 1; blki <= mt; blki++) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = 0; blkj < nt; blkj++) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_ssetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_ssetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_slarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
} // end for (Vm &Vn) > 0
} //end for blkj
} // end for blki
} //end of version 92
} // end RIGHT
magma_device_sync();
magmablasSetKernelStream( orig_stream );
magma_event_destroy( myevent[0] );
magma_event_destroy( myevent[1] );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free(dwork);
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing slarfb_gpu
*/
int main( int argc, char** argv )
{
TESTING_INIT();
// constants
const float c_zero = MAGMA_S_ZERO;
const float c_one = MAGMA_S_ONE;
const float c_neg_one = MAGMA_S_NEG_ONE;
const magma_int_t ione = 1;
// local variables
magma_int_t M, N, K, size, ldc, ldv, ldt, ldw, ldw2, nv;
magma_int_t ISEED[4] = {0,0,0,1};
float Cnorm, error, work[1];
magma_int_t status = 0;
// test all combinations of input parameters
magma_side_t side [] = { MagmaLeft, MagmaRight };
magma_trans_t trans [] = { MagmaTrans, MagmaNoTrans };
magma_direct_t direct[] = { MagmaForward, MagmaBackward };
magma_storev_t storev[] = { MagmaColumnwise, MagmaRowwise };
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% M N K storev side direct trans ||R||_F / ||HC||_F\n");
printf("%%=======================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
if ( M < K || N < K || K <= 0 ) {
printf( "%5d %5d %5d skipping because slarfb requires M >= K, N >= K, K >= 0\n",
(int) M, (int) N, (int) K );
continue;
}
for( int istor = 0; istor < 2; ++istor ) {
for( int iside = 0; iside < 2; ++iside ) {
for( int idir = 0; idir < 2; ++idir ) {
for( int itran = 0; itran < 2; ++itran ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
ldc = magma_roundup( M, opts.align ); // multiple of 32 by default
ldt = magma_roundup( K, opts.align ); // multiple of 32 by default
ldw = (side[iside] == MagmaLeft ? N : M);
ldw2 = min( M, N );
// (ldv, nv) get swapped later if rowwise
ldv = (side[iside] == MagmaLeft ? M : N);
nv = K;
// Allocate memory for matrices
float *C, *R, *V, *T, *W;
TESTING_MALLOC_CPU( C, float, ldc*N );
TESTING_MALLOC_CPU( R, float, ldc*N );
TESTING_MALLOC_CPU( V, float, ldv*K );
TESTING_MALLOC_CPU( T, float, ldt*K );
TESTING_MALLOC_CPU( W, float, ldw*K );
magmaFloat_ptr dC, dV, dT, dW, dW2;
TESTING_MALLOC_DEV( dC, float, ldc*N );
TESTING_MALLOC_DEV( dV, float, ldv*K );
TESTING_MALLOC_DEV( dT, float, ldt*K );
TESTING_MALLOC_DEV( dW, float, ldw*K );
TESTING_MALLOC_DEV( dW2, float, ldw2*K );
// C is M x N.
size = ldc*N;
lapackf77_slarnv( &ione, ISEED, &size, C );
//printf( "C=" ); magma_sprint( M, N, C, ldc );
// V is ldv x nv. See larfb docs for description.
// if column-wise and left, M x K
// if column-wise and right, N x K
// if row-wise and left, K x M
// if row-wise and right, K x N
size = ldv*nv;
lapackf77_slarnv( &ione, ISEED, &size, V );
if ( storev[istor] == MagmaColumnwise ) {
if ( direct[idir] == MagmaForward ) {
lapackf77_slaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_slaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, &V[(ldv-K)], &ldv );
}
}
else {
// rowwise, swap V's dimensions
std::swap( ldv, nv );
if ( direct[idir] == MagmaForward ) {
lapackf77_slaset( MagmaLowerStr, &K, &K, &c_zero, &c_one, V, &ldv );
}
else {
lapackf77_slaset( MagmaUpperStr, &K, &K, &c_zero, &c_one, &V[(nv-K)*ldv], &ldv );
}
}
//printf( "# ldv %d, nv %d\n", ldv, nv );
//printf( "V=" ); magma_sprint( ldv, nv, V, ldv );
// T is K x K, upper triangular for forward, and lower triangular for backward
magma_int_t k1 = K-1;
size = ldt*K;
lapackf77_slarnv( &ione, ISEED, &size, T );
if ( direct[idir] == MagmaForward ) {
lapackf77_slaset( MagmaLowerStr, &k1, &k1, &c_zero, &c_zero, &T[1], &ldt );
}
else {
lapackf77_slaset( MagmaUpperStr, &k1, &k1, &c_zero, &c_zero, &T[1*ldt], &ldt );
}
//printf( "T=" ); magma_sprint( K, K, T, ldt );
magma_ssetmatrix( M, N, C, ldc, dC, ldc, opts.queue );
magma_ssetmatrix( ldv, nv, V, ldv, dV, ldv, opts.queue );
magma_ssetmatrix( K, K, T, ldt, dT, ldt, opts.queue );
lapackf77_slarfb( lapack_side_const( side[iside] ), lapack_trans_const( trans[itran] ),
lapack_direct_const( direct[idir] ), lapack_storev_const( storev[istor] ),
&M, &N, &K,
V, &ldv, T, &ldt, C, &ldc, W, &ldw );
//printf( "HC=" ); magma_sprint( M, N, C, ldc );
if ( opts.version == 1 ) {
magma_slarfb_gpu( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw, opts.queue );
}
else {
magma_slarfb_gpu_gemm( side[iside], trans[itran], direct[idir], storev[istor],
M, N, K,
dV, ldv, dT, ldt, dC, ldc, dW, ldw, dW2, ldw2, opts.queue );
}
magma_sgetmatrix( M, N, dC, ldc, R, ldc, opts.queue );
//printf( "dHC=" ); magma_sprint( M, N, R, ldc );
// compute relative error |HC_magma - HC_lapack| / |HC_lapack|
size = ldc*N;
blasf77_saxpy( &size, &c_neg_one, C, &ione, R, &ione );
Cnorm = lapackf77_slange( "Fro", &M, &N, C, &ldc, work );
error = lapackf77_slange( "Fro", &M, &N, R, &ldc, work ) / Cnorm;
printf( "%5d %5d %5d %c %c %c %c %8.2e %s\n",
(int) M, (int) N, (int) K,
lapacke_storev_const(storev[istor]), lapacke_side_const(side[iside]),
lapacke_direct_const(direct[idir]), lapacke_trans_const(trans[itran]),
error, (error < tol ? "ok" : "failed") );
status += ! (error < tol);
TESTING_FREE_CPU( C );
TESTING_FREE_CPU( R );
TESTING_FREE_CPU( V );
TESTING_FREE_CPU( T );
TESTING_FREE_CPU( W );
TESTING_FREE_DEV( dC );
TESTING_FREE_DEV( dV );
TESTING_FREE_DEV( dT );
TESTING_FREE_DEV( dW );
TESTING_FREE_DEV( dW2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}}}}
printf( "\n" );
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}