virtual void copy(DataSpace<DIM2> &hostCurrentSize)
{
CUDA_CHECK(cudaMemcpy2DAsync(this->device->getPointer(),
this->device->getPitch(), /*this is pitch*/
this->host->getBasePointer(),
this->host->getDataSpace()[0] * sizeof (TYPE), /*this is pitch*/
hostCurrentSize[0] * sizeof (TYPE),
hostCurrentSize[1],
cudaMemcpyHostToDevice,
this->getCudaStream()));
}
virtual void copy(DataSpace<DIM2> &devCurrentSize)
{
CUDA_CHECK(cudaMemcpy2DAsync(this->destination->getPointer(),
this->destination->getPitch(),
this->source->getPointer(),
this->source->getPitch(),
devCurrentSize[0] * sizeof (TYPE),
devCurrentSize[1],
cudaMemcpyDeviceToDevice,
this->getCudaStream()));
}
void magma_copymatrix_async(
magma_int_t m, magma_int_t n, size_t elemSize,
void const* dA_src, magma_int_t lda,
void* dB_dst, magma_int_t ldb,
cudaStream_t stream )
{
cudaError_t status;
status = cudaMemcpy2DAsync(
dB_dst, ldb*elemSize,
dA_src, lda*elemSize,
m*elemSize, n, cudaMemcpyDeviceToDevice, stream );
check_error( status );
}
void magma_copymatrix_async_internal(
magma_int_t m, magma_int_t n, magma_int_t elemSize,
void const* dA_src, magma_int_t lda,
void* dB_dst, magma_int_t ldb,
cudaStream_t stream,
const char* func, const char* file, int line )
{
cudaError_t status;
status = cudaMemcpy2DAsync(
dB_dst, ldb*elemSize,
dA_src, lda*elemSize,
m*elemSize, n, cudaMemcpyDeviceToDevice, stream );
check_xerror( status, func, file, line );
}
void magma_scopymatrix_async_internal(
magma_int_t m, magma_int_t n,
float const* dA_src, magma_int_t lda,
float* dB_dst, magma_int_t ldb,
cudaStream_t stream,
const char* func, const char* file, int line )
{
cudaError_t status;
status = cudaMemcpy2DAsync(
dB_dst, ldb*sizeof(float),
dA_src, lda*sizeof(float),
m*sizeof(float), n, cudaMemcpyDeviceToDevice, stream );
check_xerror( status, func, file, line );
}
// --------------------
extern "C" void
magma_zcopymatrix_async_internal(
magma_int_t m, magma_int_t n,
magmaDoubleComplex_const_ptr dA_src, magma_int_t lda,
magmaDoubleComplex_ptr dB_dst, magma_int_t ldb,
magma_queue_t queue,
const char* func, const char* file, int line )
{
cudaError_t status;
status = cudaMemcpy2DAsync(
dB_dst, ldb*sizeof(magmaDoubleComplex),
dA_src, lda*sizeof(magmaDoubleComplex),
m*sizeof(magmaDoubleComplex), n, cudaMemcpyDeviceToDevice, queue );
check_xerror( status, func, file, line );
}
/***************************************************************************//**
@fn magma_copymatrix( m, n, elemSize, dA_src, ldda, dB_dst, lddb, queue )
Copy all or part of matrix dA_src on GPU device to dB_dst on GPU device.
Elements may be arbitrary size.
Type-safe versions set elemSize appropriately.
With CUDA unified addressing, dA and dB can be on different GPUs.
This version synchronizes the queue after the transfer.
See magma_copymatrix_async() for an asynchronous version.
@param[in]
m Number of rows of matrix A. m >= 0.
@param[in]
n Number of columns of matrix A. n >= 0.
@param[in]
elemSize Size of each element, e.g., sizeof(double).
@param[in]
dA_src Source array of dimension (ldda,n).
@param[in]
ldda Leading dimension of matrix A. ldda >= m.
@param[out]
dB_dst Destination array of dimension (lddb,n), on GPU device.
@param[in]
lddb Leading dimension of matrix B. lddb >= m.
@param[in]
queue Queue to execute in.
@ingroup magma_copymatrix
*******************************************************************************/
extern "C" void
magma_copymatrix_q_internal(
magma_int_t m, magma_int_t n, magma_int_t elemSize,
magma_const_ptr dA_src, magma_int_t ldda,
magma_ptr dB_dst, magma_int_t lddb,
magma_queue_t queue,
const char* func, const char* file, int line )
{
assert( queue != NULL );
cudaError_t status;
status = cudaMemcpy2DAsync(
dB_dst, int(lddb*elemSize),
dA_src, int(ldda*elemSize),
int(m*elemSize), int(n), cudaMemcpyDeviceToDevice, queue->cuda_stream() );
cudaStreamSynchronize( queue->cuda_stream() );
check_xerror( status, func, file, line );
}
SEXP R_auto_cudaMemcpy2DAsync(SEXP r_dst, SEXP r_dpitch, SEXP r_src, SEXP r_spitch, SEXP r_width, SEXP r_height, SEXP r_kind, SEXP r_stream)
{
SEXP r_ans = R_NilValue;
void * dst = GET_REF(r_dst, void );
size_t dpitch = REAL(r_dpitch)[0];
const void * src = GET_REF(r_src, const void );
size_t spitch = REAL(r_spitch)[0];
size_t width = REAL(r_width)[0];
size_t height = REAL(r_height)[0];
enum cudaMemcpyKind kind = (enum cudaMemcpyKind) INTEGER(r_kind)[0];
cudaStream_t stream = (cudaStream_t) getRReference(r_stream);
cudaError_t ans;
ans = cudaMemcpy2DAsync(dst, dpitch, src, spitch, width, height, kind, stream);
r_ans = Renum_convert_cudaError_t(ans) ;
return(r_ans);
}
/***************************************************************************//**
@fn magma_copymatrix_async( m, n, elemSize, dA_src, ldda, dB_dst, lddb, queue )
Copy all or part of matrix dA_src on GPU device to dB_dst on GPU device.
Elements may be arbitrary size.
Type-safe versions set elemSize appropriately.
With CUDA unified addressing, dA and dB can be on different GPUs.
This version is asynchronous: it may return before the transfer finishes.
See magma_copyvector() for a synchronous version.
@param[in]
m Number of rows of matrix A. m >= 0.
@param[in]
n Number of columns of matrix A. n >= 0.
@param[in]
elemSize Size of each element, e.g., sizeof(double).
@param[in]
dA_src Source array of dimension (ldda,n), on GPU device.
@param[in]
ldda Leading dimension of matrix A. ldda >= m.
@param[out]
dB_dst Destination array of dimension (lddb,n), on GPU device.
@param[in]
lddb Leading dimension of matrix B. lddb >= m.
@param[in]
queue Queue to execute in.
@ingroup magma_copymatrix
*******************************************************************************/
extern "C" void
magma_copymatrix_async_internal(
magma_int_t m, magma_int_t n, magma_int_t elemSize,
magma_const_ptr dA_src, magma_int_t ldda,
magma_ptr dB_dst, magma_int_t lddb,
magma_queue_t queue,
const char* func, const char* file, int line )
{
// for backwards compatability, accepts NULL queue to mean NULL stream.
cudaStream_t stream = NULL;
if ( queue != NULL ) {
stream = queue->cuda_stream();
}
else {
fprintf( stderr, "Warning: %s got NULL queue\n", __func__ );
}
cudaError_t status;
status = cudaMemcpy2DAsync(
dB_dst, int(lddb*elemSize),
dA_src, int(ldda*elemSize),
int(m*elemSize), int(n), cudaMemcpyDeviceToDevice, stream );
check_xerror( status, func, file, line );
}
void magmablas_ssymm_mgpu_spec(
magma_side_t side, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
float alpha,
float *dA[], magma_int_t ldda, magma_int_t offset,
float *dB[], magma_int_t lddb,
float beta,
float *dC[], magma_int_t lddc,
float *dwork[], magma_int_t dworksiz,
float *C, magma_int_t ldc,
float *work[], magma_int_t ldwork,
magma_int_t ngpu, magma_int_t nb,
magma_queue_t streams[][20], magma_int_t nstream,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10],magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
if ( side != MagmaLeft || uplo != MagmaLower ) {
fprintf( stderr, "%s: only Left Lower implemented\n", __func__ );
}
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nstream >= ngpu );
assert( nbevents >= ngpu*ngpu );
float *dwork1[MagmaMaxGPUs];
float *dwork2[MagmaMaxGPUs];
magma_int_t lddwork = lddc;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev];
dwork2[dev] = dwork[dev]+n*lddwork;
}
assert( dworksiz >= (2*n*lddwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev,devperm,myblk,mycolsize,myblkoffst;
magma_int_t gdev,gcolsize,gmaster,gngpu;
magma_int_t masterdev,lcdev,lccolsize,myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m,nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset),nb);
magma_int_t maxgsize = n*nb*magma_ceildiv(nbblk,ngpu);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
//*******************************
// each GPU make a GEMM with the
// transpose of its blocks to compute
// a final portion of X=A*VT
//*******************************
/* dB = V*T already ==> dB' = T'*V'
* compute T'*V'*X is equal to compute locally (VT)'_i*X_i
* then each GPU broadcast its X_i to assemble the full X which is used
* to compute W = X - 0.5 * V * T'*V'*X = X - 0.5 * V *dwork3
*/
if(ngpu ==1){
magma_setdevice( 0 );
magmablasSetKernelStream( streams[ 0 ][ 0 ] );
// compute X[me] = A*VT = A[me]^tr *VT;
magma_sgemm( MagmaTrans, MagmaNoTrans, m, n, m,
alpha, dA(0,offset,offset), ldda,
dB[0], lddb,
beta, dC[0], lddc );
return;
}
//ngpu>1
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize,m);
if(mycolsize>0){
if(masterdev==-1) masterdev = dev;
//printf("dev %d devperm %d on cmplx %d master %d nbblk %d myblk %d m %d n %d mycolsize %d stdev %d fstblksize %d lastdev %d lastsize %d dA(%d,%d,%d) ==> dwork(%d,%d)\n",dev,devperm,cmplxid,masterdev,nbblk,myblk,m,n,mycolsize,stdev,fstblksiz,devlstblk,remm%nb,dev,offset,myblkoffst,dev,maxgsize*dev);
gpuisactive[dev] = mycolsize;
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ dev ] );
magma_sgemm( MagmaTrans, MagmaNoTrans, mycolsize, n, m,
alpha, dA(dev,offset,myblkoffst), ldda,
dB(dev,0,0), lddb,
beta, &dwork[dev][maxgsize*dev], mycolsize );
magma_event_record(redevents[dev][dev*ngpu+dev], streams[dev][dev]);
}
if(dev == masterdev){
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[ dev ][ dev ] );
}
*/
//*******************************
// each Master GPU has the final
// result either by receiving
// from CPU of by making the add
// by himself, so now it is time
// to broadcast over the GPUs of
// its board.
//*******************************
//printf("=======================================================================\n");
//printf(" sending \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU send my portion local
// to all active gpu of my cmplex and global to the
// active master of the other real and they should
// send it out to their actives slaves.
magma_setdevice( dev );
//==============================================
// sending to the master of the active real
//==============================================
//printf ("\n\n**************GPU %d\n ",dev);
//printf (" GPU %d sending to cmplx masters\n",dev);
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//printf (" device %d from cmplx %d is sending to master %d on cmplx %d block of size %d event %d\n",dev,cmplxid,gmaster,k,mycolsize,redevents[dev][gmaster*ngpu+dev]);
magma_queue_wait_event(streams[ dev ][ gmaster ], redevents[dev][dev*ngpu+dev]);
cudaMemcpy2DAsync(&dwork[gmaster][maxgsize*dev], mycolsize*sizeof(float),
&dwork[dev][maxgsize*dev], mycolsize*sizeof(float),
mycolsize*sizeof(float), n,
cudaMemcpyDeviceToDevice, streams[dev][gmaster]);
magma_event_record(redevents[dev][gmaster*ngpu+dev], streams[dev][gmaster]);
}
}
}
//==============================================
//
//==============================================
// sending to the active GPUs of my real
//==============================================
//printf (" GPU %d sending internal\n",dev);
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=dev)&&(lccolsize>0)){
//printf (" device %d from cmplx %d is sending internal to dev %d block of size %d event %d\n",dev,cmplxid,lcdev,mycolsize,redevents[dev][lcdev*ngpu+dev]);
magma_queue_wait_event(streams[ dev ][ lcdev ], redevents[dev][dev*ngpu+dev]);
cudaMemcpy2DAsync(&dwork[lcdev][maxgsize*dev], mycolsize*sizeof(float),
&dwork[dev][maxgsize*dev], mycolsize*sizeof(float),
mycolsize*sizeof(float), n,
cudaMemcpyDeviceToDevice, streams[dev][lcdev]);
magma_event_record(redevents[dev][lcdev*ngpu+dev], streams[dev][lcdev]);
}
}
//==============================================
}// end if mycolsize>0
}// for idev
}// for cmplxid
//printf("=======================================================================\n");
//printf(" master wait and resend internally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//==============================================
// if I am active master so wait receiving contribution
// of the GPUs of other real and send it locally
//==============================================
if(masterdev != -1){
mycolsize = gpuisactive[masterdev];
magma_setdevice( masterdev );
//printf(" GPU %d distributing internal\n",masterdev);
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gngpu = gnode[k][MagmaMaxGPUs];
for( magma_int_t g = 0; g < gngpu; ++g ) {
gdev = gnode[k][g];
gcolsize = gpuisactive[gdev];
// check if I received from this GPU,
// if yes send it to my group
if(gcolsize>0){
magma_queue_wait_event(streams[ masterdev ][ gdev ], redevents[gdev][masterdev*ngpu+gdev]);
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" Master %d on cmplx %d waiting on event %d is distributing internal results of %d to lcdev %d block of size %d event %d\n", masterdev,cmplxid,redevents[gdev][masterdev*ngpu+gdev],gdev,lcdev,gcolsize,redevents[masterdev][lcdev*ngpu+gdev]);
cudaMemcpy2DAsync(&dwork[lcdev][maxgsize*gdev], gcolsize*sizeof(float),
&dwork[masterdev][maxgsize*gdev], gcolsize*sizeof(float),
gcolsize*sizeof(float), n,
cudaMemcpyDeviceToDevice, streams[masterdev][gdev]);
magma_event_record(redevents[masterdev][lcdev*ngpu+gdev], streams[masterdev][gdev]);
}
}
}
}
}
}
}// if active master
//==============================================
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[ dev ][ 0 ] );
for( magma_int_t s = 0; s < ngpu; ++s ) {
magma_queue_sync( streams[ dev ][ s ] );
}
}
*/
//printf("=======================================================================\n");
//printf(" distributing \n");
//printf("=======================================================================\n");
magma_int_t lcblki,gbblki,gblk,ib;
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){ // I am an active GPU
//printf("\n\n==============GPU %d collecting\n",dev);
magma_setdevice( dev );
// collect my results first as tyhere is no need to wait to
// receive nothing, just wait that my gemm are done.
// in theory this should be inside the loop but cuda was not
// able to run it first for all gpu and on gpu>0 it was waiting
// however it was on different stream so it should run. but maybe
// this is because there are too many function call and this make
// cuda not handleit so nice. anyway it coul dbe removed when cuda
// is able to lunch it first without wait.
gdev = dev;
gcolsize = gpuisactive[gdev];
if(gcolsize>0){
devperm = (gdev-stdev+ngpu)%ngpu;
gblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
magmablasSetKernelStream( streams[ dev ][ gdev ] );
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
//printf (" GPU %d stream %d doing slacpy\n",dev,streams[ dev ][ gdev ]);
for( magma_int_t blki = 0; blki < gblk; ++blki){
gbblki = (blki*ngpu + devperm)*nb - blockoffset;
lcblki = blki*nb;
ib = nb;//min(nb,m-gbblki);
if(gdev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
ib = min(ib,m-gbblki);
//printf(" blockoffset %d nbblk %d stdev %d receiving from gdev %d gblk %d gcolsize %d copying blki %d of size ibxn %dx%d from work[%d] to C[%d]\n", blockoffset,nbblk,stdev,gdev,gblk,gcolsize,blki,ib,n,lcblki,gbblki);
magmablas_slacpy( MagmaFull, ib, n, &dwork[dev][maxgsize*gdev+lcblki], gcolsize, &dC[dev][gbblki], lddc);
}// end blki
}
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
gngpu = gnode[k][MagmaMaxGPUs];
for( magma_int_t g = 0; g < gngpu; ++g ) {
gdev = gnode[k][g];
gcolsize = gpuisactive[gdev];
// if gcolsize>0, ==> gpu gdev was active and so
// I received from him/computed a portion of dwork,
// so go over its gblk and distribute it on dC.
if(gdev!=dev){
if(gcolsize>0){
devperm = (gdev-stdev+ngpu)%ngpu;
gblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
magmablasSetKernelStream( streams[ dev ][ gdev ] );
if(k==cmplxid){
//we are on the same group so wait on event issued by gdev for me citing his id
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
//printf (" GPU %d stream %d waiting on event %d to collecte from %d the size of gcolsize %d\n",dev,streams[ dev ][ gdev ],redevents[gdev][dev*ngpu+gdev],gdev,gcolsize);
}else{
//we are on different group so:
//if I am the master wait on the event issued by gdev for me citing his id
//else wait event issued by my master for me on the behalf of gdev
//printf (" GPU %d stream %d waiting on event %d to collecte from %d the size of gcolsize %d\n",dev,streams[ dev ][ gdev ],redevents[masterdev][dev*ngpu+gdev],gdev,gcolsize);
if(dev==masterdev)
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
else
magma_queue_wait_event(streams[ dev ][ gdev ], redevents[masterdev][dev*ngpu+gdev]);
}
//printf (" GPU %d stream %d doing slacpy\n",dev,streams[ dev ][ gdev ]);
for( magma_int_t blki = 0; blki < gblk; ++blki){
gbblki = (blki*ngpu + devperm)*nb - blockoffset;
lcblki = blki*nb;
ib = nb;//min(nb,m-gbblki);
if(gdev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
ib = min(ib,m-gbblki);
//printf(" blockoffset %d nbblk %d stdev %d receiving from gdev %d gblk %d gcolsize %d copying blki %d of size ibxn %dx%d from work[%d] to C[%d]\n", blockoffset,nbblk,stdev,gdev,gblk,gcolsize,blki,ib,n,lcblki,gbblki);
magmablas_slacpy( MagmaFull, ib, n, &dwork[dev][maxgsize*gdev+lcblki], gcolsize, &dC[dev][gbblki], lddc);
}// end blki
}// en gcolsize>0 meaning gdev is active
} // end if gdev != dev
}// end loop over the g gpus of the cmplx k
}//end loop over the real k
}// end mycolsize>0 meaning that I am active
}// end loop over idev of cmplxid
}// end loop of the cmplx
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
cudaDeviceSynchronize();
}
// put back the input gpu and its input stream
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}
extern "C" magma_int_t
magma_zgeqrf2(magma_context *cntxt, magma_int_t m, magma_int_t n,
cuDoubleComplex *a, magma_int_t lda, cuDoubleComplex *tau,
cuDoubleComplex *work, magma_int_t lwork,
magma_int_t *info)
{
/* -- MAGMA (version 1.5.0-beta3) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date July 2014
Purpose
=======
ZGEQRF computes a QR factorization of a COMPLEX_16 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.
Arguments
=========
CNTXT (input) MAGMA_CONTEXT
CNTXT specifies the MAGMA hardware context for this routine.
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX_16 array, 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 cudaMallocHost.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
TAU (output) COMPLEX_16 array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX_16 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 cudaMallocHost.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= N*NB,
where NB can be obtained through magma_get_zgeqrf_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
if INFO = -8, the GPU 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))
int cnt=-1;
cuDoubleComplex c_one = MAGMA_Z_ONE;
int i, k, lddwork, old_i, old_ib;
int nbmin, nx, ib, ldda;
*info = 0;
magma_qr_params *qr_params = (magma_qr_params *)cntxt->params;
int nb = qr_params->nb;
int lwkopt = n * nb;
work[0] = MAGMA_Z_MAKE( (double)lwkopt, 0 );
long 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 MAGMA_ERR_ILLEGAL_VALUE;
}
else if (lquery)
return MAGMA_SUCCESS;
k = min(m,n);
if (k == 0) {
work[0] = c_one;
return MAGMA_SUCCESS;
}
cublasStatus status;
static cudaStream_t stream[2];
cudaStreamCreate(&stream[0]);
cudaStreamCreate(&stream[1]);
nbmin = 2;
nx = nb;
lddwork = ((n+31)/32)*32;
ldda = ((m+31)/32)*32;
cuDoubleComplex *da;
status = cublasAlloc((n)*ldda + nb*lddwork, sizeof(cuDoubleComplex), (void**)&da);
if (status != CUBLAS_STATUS_SUCCESS) {
*info = -8;
return 0;
}
cuDoubleComplex *dwork = da + ldda*(n);
if (nb >= nbmin && nb < k && nx < k) {
/* Use blocked code initially */
cudaMemcpy2DAsync(da_ref(0,nb), ldda*sizeof(cuDoubleComplex),
a_ref(0,nb), lda *sizeof(cuDoubleComplex),
sizeof(cuDoubleComplex)*(m), (n-nb),
cudaMemcpyHostToDevice,stream[0]);
old_i = 0; old_ib = nb;
for (i = 0; i < k-nx; i += nb) {
ib = min(k-i, nb);
if (i>0){
cudaMemcpy2DAsync( a_ref(i,i), lda *sizeof(cuDoubleComplex),
da_ref(i,i), ldda*sizeof(cuDoubleComplex),
sizeof(cuDoubleComplex)*(m-i), ib,
cudaMemcpyDeviceToHost,stream[1]);
cudaMemcpy2DAsync( a_ref(0,i), lda *sizeof(cuDoubleComplex),
da_ref(0,i), ldda*sizeof(cuDoubleComplex),
sizeof(cuDoubleComplex)*i, ib,
cudaMemcpyDeviceToHost,stream[0]);
/* Apply H' to A(i:m,i+2*ib:n) from the left */
magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-old_i, n-old_i-2*old_ib, old_ib,
da_ref(old_i, old_i), ldda, dwork, lddwork,
da_ref(old_i, old_i+2*old_ib), ldda, dwork+old_ib, lddwork);
}
cudaStreamSynchronize(stream[1]);
int rows = m-i;
cnt++;
cntxt->nb = qr_params->ib;
magma_zgeqrf_mc(cntxt, &rows, &ib, a_ref(i,i), &lda,
tau+i, work, &lwork, info);
cntxt->nb = nb;
/* Form the triangular factor of the block reflector
H = H(i) H(i+1) . . . H(i+ib-1) */
lapackf77_zlarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib, a_ref(i,i), &lda, tau+i, qr_params->t+cnt*nb*nb, &ib);
if (cnt < qr_params->np_gpu) {
qr_params->p[cnt]=a;
}
zpanel_to_q(MagmaUpper, ib, a_ref(i,i), lda, qr_params->w+cnt*qr_params->nb*qr_params->nb);
cublasSetMatrix(rows, ib, sizeof(cuDoubleComplex),
a_ref(i,i), lda, da_ref(i,i), ldda);
if (qr_params->flag == 1)
zq_to_panel(MagmaUpper, ib, a_ref(i,i), lda, qr_params->w+cnt*qr_params->nb*qr_params->nb);
if (i + ib < n) {
cublasSetMatrix(ib, ib, sizeof(cuDoubleComplex), qr_params->t+cnt*nb*nb, ib, dwork, lddwork);
if (i+ib < k-nx)
/* Apply H' to A(i:m,i+ib:i+2*ib) from the left */
magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, ib, ib,
da_ref(i, i ), ldda, dwork, lddwork,
da_ref(i, i+ib), ldda, dwork+ib, lddwork);
else
magma_zlarfb_gpu( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
rows, n-i-ib, ib,
da_ref(i, i ), ldda, dwork, lddwork,
da_ref(i, i+ib), ldda, dwork+ib, lddwork);
old_i = i;
old_ib = ib;
}
}
} else {
i = 0;
}
/* Use unblocked code to factor the last or only block. */
if (i < k)
{
ib = n-i;
if (i!=0)
cublasGetMatrix(m, ib, sizeof(cuDoubleComplex),
da_ref(0,i), ldda, a_ref(0,i), lda);
int rows = m-i;
cnt++;
lapackf77_zgeqrf(&rows, &ib, a_ref(i,i), &lda, tau+i, work, &lwork, info);
if (cnt < qr_params->np_gpu)
{
int ib2=min(ib,nb);
lapackf77_zlarft( MagmaForwardStr, MagmaColumnwiseStr,
&rows, &ib2, a_ref(i,i), &lda, tau+i, qr_params->t+cnt*nb*nb, &ib2);
qr_params->p[cnt]=a;
}
}
cudaStreamDestroy( stream[0] );
cudaStreamDestroy( stream[1] );
cublasFree(da);
return MAGMA_SUCCESS;
} /* magma_zgeqrf */
int TEMPLATE2 (CHOLMOD (gpu_lower_potrf))
(
Int nscol2, /* S is nscol2-by-nscol2 */
Int nsrow, /* leading dimension of S */
Int psx, /* S is located at Lx + L_ENTRY*psx */
double *Lx, /* contains S; overwritten with Cholesky factor */
Int *info, /* BLAS info return value */
cholmod_common *Common,
cholmod_gpu_pointers *gpu_p
)
{
double *devPtrA, *devPtrB, *A ;
double alpha, beta ;
cudaError_t cudaStat ;
cublasStatus_t cublasStatus ;
Int j, nsrow2, nb, n, gpu_lda, lda, gpu_ldb ;
int ilda, ijb, iinfo ;
#ifndef NTIMER
double tstart ;
#endif
if (nscol2 * L_ENTRY < CHOLMOD_POTRF_LIMIT)
{
/* too small for the CUDA BLAS; use the CPU instead */
return (0) ;
}
#ifndef NTIMER
tstart = SuiteSparse_time ( ) ;
Common->CHOLMOD_GPU_POTRF_CALLS++ ;
#endif
nsrow2 = nsrow - nscol2 ;
/* ---------------------------------------------------------------------- */
/* heuristic to get the block size depending of the problem size */
/* ---------------------------------------------------------------------- */
nb = 128 ;
if (nscol2 > 4096) nb = 256 ;
if (nscol2 > 8192) nb = 384 ;
n = nscol2 ;
gpu_lda = ((nscol2+31)/32)*32 ;
lda = nsrow ;
A = gpu_p->h_Lx[(Common->ibuffer+CHOLMOD_HOST_SUPERNODE_BUFFERS-1)%
CHOLMOD_HOST_SUPERNODE_BUFFERS];
/* ---------------------------------------------------------------------- */
/* determine the GPU leading dimension of B */
/* ---------------------------------------------------------------------- */
gpu_ldb = 0 ;
if (nsrow2 > 0)
{
gpu_ldb = ((nsrow2+31)/32)*32 ;
}
/* ---------------------------------------------------------------------- */
/* remember where device memory is, to be used by triangular solve later */
/* ---------------------------------------------------------------------- */
devPtrA = gpu_p->d_Lx[0];
devPtrB = gpu_p->d_Lx[1];
/* ---------------------------------------------------------------------- */
/* copy A from device to device */
/* ---------------------------------------------------------------------- */
cudaStat = cudaMemcpy2DAsync ( devPtrA,
gpu_lda * L_ENTRY * sizeof (devPtrA[0]),
gpu_p->d_A[1],
nsrow * L_ENTRY * sizeof (Lx[0]),
nscol2 * L_ENTRY * sizeof (devPtrA[0]),
nscol2,
cudaMemcpyDeviceToDevice,
Common->gpuStream[0] );
if ( cudaStat ) {
ERROR ( CHOLMOD_GPU_PROBLEM, "GPU memcopy device to device");
}
/* ---------------------------------------------------------------------- */
/* copy B in advance, for gpu_triangular_solve */
/* ---------------------------------------------------------------------- */
if (nsrow2 > 0)
{
cudaStat = cudaMemcpy2DAsync (devPtrB,
gpu_ldb * L_ENTRY * sizeof (devPtrB [0]),
gpu_p->d_A[1] + L_ENTRY*nscol2,
nsrow * L_ENTRY * sizeof (Lx [0]),
nsrow2 * L_ENTRY * sizeof (devPtrB [0]),
nscol2,
cudaMemcpyDeviceToDevice,
Common->gpuStream[0]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU memcopy to device") ;
}
}
/* ------------------------------------------------------------------ */
/* define the dpotrf stream */
/* ------------------------------------------------------------------ */
cublasStatus = cublasSetStream (Common->cublasHandle,
Common->gpuStream [0]) ;
if (cublasStatus != CUBLAS_STATUS_SUCCESS) {
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS stream") ;
}
/* ---------------------------------------------------------------------- */
/* block Cholesky factorization of S */
/* ---------------------------------------------------------------------- */
for (j = 0 ; j < n ; j += nb)
{
Int jb = nb < (n-j) ? nb : (n-j) ;
/* ------------------------------------------------------------------ */
/* do the CUDA BLAS dsyrk */
/* ------------------------------------------------------------------ */
alpha = -1.0 ;
beta = 1.0 ;
#ifdef REAL
cublasStatus = cublasDsyrk (Common->cublasHandle,
CUBLAS_FILL_MODE_LOWER, CUBLAS_OP_N, jb, j,
&alpha, devPtrA + j, gpu_lda,
&beta, devPtrA + j + j*gpu_lda, gpu_lda) ;
#else
cublasStatus = cublasZherk (Common->cublasHandle,
CUBLAS_FILL_MODE_LOWER, CUBLAS_OP_N, jb, j,
&alpha, (cuDoubleComplex*)devPtrA + j,
gpu_lda,
&beta,
(cuDoubleComplex*)devPtrA + j + j*gpu_lda,
gpu_lda) ;
#endif
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS routine failure") ;
}
/* ------------------------------------------------------------------ */
cudaStat = cudaEventRecord (Common->cublasEventPotrf [0],
Common->gpuStream [0]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA event failure") ;
}
cudaStat = cudaStreamWaitEvent (Common->gpuStream [1],
Common->cublasEventPotrf [0], 0) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA event failure") ;
}
/* ------------------------------------------------------------------ */
/* copy back the jb columns on two different streams */
/* ------------------------------------------------------------------ */
cudaStat = cudaMemcpy2DAsync (A + L_ENTRY*(j + j*lda),
lda * L_ENTRY * sizeof (double),
devPtrA + L_ENTRY*(j + j*gpu_lda),
gpu_lda * L_ENTRY * sizeof (double),
L_ENTRY * sizeof (double)*jb,
jb,
cudaMemcpyDeviceToHost,
Common->gpuStream [1]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU memcopy from device") ;
}
/* ------------------------------------------------------------------ */
/* do the CUDA BLAS dgemm */
/* ------------------------------------------------------------------ */
if ((j+jb) < n)
{
#ifdef REAL
alpha = -1.0 ;
beta = 1.0 ;
cublasStatus = cublasDgemm (Common->cublasHandle,
CUBLAS_OP_N, CUBLAS_OP_T,
(n-j-jb), jb, j,
&alpha,
devPtrA + (j+jb), gpu_lda,
devPtrA + (j) , gpu_lda,
&beta,
devPtrA + (j+jb + j*gpu_lda), gpu_lda) ;
#else
cuDoubleComplex calpha = {-1.0,0.0} ;
cuDoubleComplex cbeta = { 1.0,0.0} ;
cublasStatus = cublasZgemm (Common->cublasHandle,
CUBLAS_OP_N, CUBLAS_OP_C,
(n-j-jb), jb, j,
&calpha,
(cuDoubleComplex*)devPtrA + (j+jb),
gpu_lda,
(cuDoubleComplex*)devPtrA + (j),
gpu_lda,
&cbeta,
(cuDoubleComplex*)devPtrA +
(j+jb + j*gpu_lda),
gpu_lda ) ;
#endif
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS routine failure") ;
}
}
cudaStat = cudaStreamSynchronize (Common->gpuStream [1]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU memcopy to device") ;
}
/* ------------------------------------------------------------------ */
/* compute the Cholesky factorization of the jbxjb block on the CPU */
/* ------------------------------------------------------------------ */
ilda = (int) lda ;
ijb = jb ;
#ifdef REAL
LAPACK_DPOTRF ("L", &ijb, A + L_ENTRY * (j + j*lda), &ilda, &iinfo) ;
#else
LAPACK_ZPOTRF ("L", &ijb, A + L_ENTRY * (j + j*lda), &ilda, &iinfo) ;
#endif
*info = iinfo ;
if (*info != 0)
{
*info = *info + j ;
break ;
}
/* ------------------------------------------------------------------ */
/* copy the result back to the GPU */
/* ------------------------------------------------------------------ */
cudaStat = cudaMemcpy2DAsync (devPtrA + L_ENTRY*(j + j*gpu_lda),
gpu_lda * L_ENTRY * sizeof (double),
A + L_ENTRY * (j + j*lda),
lda * L_ENTRY * sizeof (double),
L_ENTRY * sizeof (double) * jb,
jb,
cudaMemcpyHostToDevice,
Common->gpuStream [0]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU memcopy to device") ;
}
/* ------------------------------------------------------------------ */
/* do the CUDA BLAS dtrsm */
/* ------------------------------------------------------------------ */
if ((j+jb) < n)
{
#ifdef REAL
alpha = 1.0 ;
cublasStatus = cublasDtrsm (Common->cublasHandle,
CUBLAS_SIDE_RIGHT,
CUBLAS_FILL_MODE_LOWER,
CUBLAS_OP_T, CUBLAS_DIAG_NON_UNIT,
(n-j-jb), jb,
&alpha,
devPtrA + (j + j*gpu_lda), gpu_lda,
devPtrA + (j+jb + j*gpu_lda), gpu_lda) ;
#else
cuDoubleComplex calpha = {1.0,0.0};
cublasStatus = cublasZtrsm (Common->cublasHandle,
CUBLAS_SIDE_RIGHT,
CUBLAS_FILL_MODE_LOWER,
CUBLAS_OP_C, CUBLAS_DIAG_NON_UNIT,
(n-j-jb), jb,
&calpha,
(cuDoubleComplex *)devPtrA +
(j + j*gpu_lda),
gpu_lda,
(cuDoubleComplex *)devPtrA +
(j+jb + j*gpu_lda),
gpu_lda) ;
#endif
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS routine failure") ;
}
/* -------------------------------------------------------------- */
/* Copy factored column back to host. */
/* -------------------------------------------------------------- */
cudaStat = cudaEventRecord (Common->cublasEventPotrf[2],
Common->gpuStream[0]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA event failure") ;
}
cudaStat = cudaStreamWaitEvent (Common->gpuStream[1],
Common->cublasEventPotrf[2], 0) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA event failure") ;
}
cudaStat = cudaMemcpy2DAsync (A + L_ENTRY*(j + jb + j * lda),
lda * L_ENTRY * sizeof (double),
devPtrA + L_ENTRY*
(j + jb + j * gpu_lda),
gpu_lda * L_ENTRY * sizeof (double),
L_ENTRY * sizeof (double)*
(n - j - jb), jb,
cudaMemcpyDeviceToHost,
Common->gpuStream[1]) ;
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU memcopy to device") ;
}
}
}
#ifndef NTIMER
Common->CHOLMOD_GPU_POTRF_TIME += SuiteSparse_time ( ) - tstart ;
#endif
return (1) ;
}
int TEMPLATE2 (CHOLMOD (gpu_triangular_solve))
(
Int nsrow2, /* L1 and S2 are nsrow2-by-nscol2 */
Int nscol2, /* L1 is nscol2-by-nscol2 */
Int nsrow, /* leading dimension of L1, L2, and S2 */
Int psx, /* L1 is at Lx+L_ENTRY*psx;
* L2 at Lx+L_ENTRY*(psx+nscol2)*/
double *Lx, /* holds L1, L2, and S2 */
cholmod_common *Common,
cholmod_gpu_pointers *gpu_p
)
{
double *devPtrA, *devPtrB ;
cudaError_t cudaStat ;
cublasStatus_t cublasStatus ;
Int gpu_lda, gpu_ldb, gpu_rowstep ;
Int gpu_row_start = 0 ;
Int gpu_row_max_chunk, gpu_row_chunk;
int ibuf = 0;
int iblock = 0;
int iHostBuff = (Common->ibuffer+CHOLMOD_HOST_SUPERNODE_BUFFERS-1) %
CHOLMOD_HOST_SUPERNODE_BUFFERS;
int i, j;
Int iidx;
int iwrap;
#ifndef NTIMER
double tstart ;
#endif
#ifdef REAL
double alpha = 1.0 ;
gpu_row_max_chunk = 768;
#else
cuDoubleComplex calpha = {1.0,0.0} ;
gpu_row_max_chunk = 256;
#endif
if ( nsrow2 <= 0 )
{
return (0) ;
}
#ifndef NTIMER
tstart = SuiteSparse_time ( ) ;
Common->CHOLMOD_GPU_TRSM_CALLS++ ;
#endif
gpu_lda = ((nscol2+31)/32)*32 ;
gpu_ldb = ((nsrow2+31)/32)*32 ;
devPtrA = gpu_p->d_Lx[0];
devPtrB = gpu_p->d_Lx[1];
/* make sure the copy of B has completed */
cudaStreamSynchronize( Common->gpuStream[0] );
/* ---------------------------------------------------------------------- */
/* do the CUDA BLAS dtrsm */
/* ---------------------------------------------------------------------- */
while ( gpu_row_start < nsrow2 )
{
gpu_row_chunk = nsrow2 - gpu_row_start;
if ( gpu_row_chunk > gpu_row_max_chunk ) {
gpu_row_chunk = gpu_row_max_chunk;
}
cublasStatus = cublasSetStream ( Common->cublasHandle,
Common->gpuStream[ibuf] );
if ( cublasStatus != CUBLAS_STATUS_SUCCESS )
{
ERROR ( CHOLMOD_GPU_PROBLEM, "GPU CUBLAS stream");
}
#ifdef REAL
cublasStatus = cublasDtrsm (Common->cublasHandle,
CUBLAS_SIDE_RIGHT,
CUBLAS_FILL_MODE_LOWER,
CUBLAS_OP_T,
CUBLAS_DIAG_NON_UNIT,
gpu_row_chunk,
nscol2,
&alpha,
devPtrA,
gpu_lda,
devPtrB + gpu_row_start,
gpu_ldb) ;
#else
cublasStatus = cublasZtrsm (Common->cublasHandle,
CUBLAS_SIDE_RIGHT,
CUBLAS_FILL_MODE_LOWER,
CUBLAS_OP_C,
CUBLAS_DIAG_NON_UNIT,
gpu_row_chunk,
nscol2,
&calpha,
(const cuDoubleComplex *) devPtrA,
gpu_lda,
(cuDoubleComplex *)devPtrB + gpu_row_start ,
gpu_ldb) ;
#endif
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS routine failure") ;
}
/* ------------------------------------------------------------------ */
/* copy result back to the CPU */
/* ------------------------------------------------------------------ */
cudaStat = cudaMemcpy2DAsync (
gpu_p->h_Lx[iHostBuff] +
L_ENTRY*(nscol2+gpu_row_start),
nsrow * L_ENTRY * sizeof (Lx [0]),
devPtrB + L_ENTRY*gpu_row_start,
gpu_ldb * L_ENTRY * sizeof (devPtrB [0]),
gpu_row_chunk * L_ENTRY *
sizeof (devPtrB [0]),
nscol2,
cudaMemcpyDeviceToHost,
Common->gpuStream[ibuf]);
if (cudaStat)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU memcopy from device") ;
}
cudaEventRecord ( Common->updateCBuffersFree[ibuf],
Common->gpuStream[ibuf] );
gpu_row_start += gpu_row_chunk;
ibuf++;
ibuf = ibuf % CHOLMOD_HOST_SUPERNODE_BUFFERS;
iblock ++;
if ( iblock >= CHOLMOD_HOST_SUPERNODE_BUFFERS )
{
Int gpu_row_start2 ;
Int gpu_row_end ;
/* then CHOLMOD_HOST_SUPERNODE_BUFFERS worth of work has been
* scheduled, so check for completed events and copy result into
* Lx before continuing. */
cudaEventSynchronize ( Common->updateCBuffersFree
[iblock%CHOLMOD_HOST_SUPERNODE_BUFFERS] );
/* copy into Lx */
gpu_row_start2 = nscol2 +
(iblock-CHOLMOD_HOST_SUPERNODE_BUFFERS)
*gpu_row_max_chunk;
gpu_row_end = gpu_row_start2+gpu_row_max_chunk;
if ( gpu_row_end > nsrow ) gpu_row_end = nsrow;
#pragma omp parallel for num_threads(CHOLMOD_OMP_NUM_THREADS) \
private(iidx) if ( nscol2 > 32 )
for ( j=0; j<nscol2; j++ ) {
for ( i=gpu_row_start2*L_ENTRY; i<gpu_row_end*L_ENTRY; i++ ) {
iidx = j*nsrow*L_ENTRY+i;
Lx[psx*L_ENTRY+iidx] = gpu_p->h_Lx[iHostBuff][iidx];
}
}
}
}
/* Convenient to copy the L1 block here */
#pragma omp parallel for num_threads(CHOLMOD_OMP_NUM_THREADS) \
private ( iidx ) if ( nscol2 > 32 )
for ( j=0; j<nscol2; j++ ) {
for ( i=j*L_ENTRY; i<nscol2*L_ENTRY; i++ ) {
iidx = j*nsrow*L_ENTRY + i;
Lx[psx*L_ENTRY+iidx] = gpu_p->h_Lx[iHostBuff][iidx];
}
}
/* now account for the last HSTREAMS buffers */
for ( iwrap=0; iwrap<CHOLMOD_HOST_SUPERNODE_BUFFERS; iwrap++ )
{
int i, j;
Int gpu_row_start2 = nscol2 + (iblock-CHOLMOD_HOST_SUPERNODE_BUFFERS)
*gpu_row_max_chunk;
if (iblock-CHOLMOD_HOST_SUPERNODE_BUFFERS >= 0 &&
gpu_row_start2 < nsrow )
{
Int iidx;
Int gpu_row_end = gpu_row_start2+gpu_row_max_chunk;
if ( gpu_row_end > nsrow ) gpu_row_end = nsrow;
cudaEventSynchronize ( Common->updateCBuffersFree
[iblock%CHOLMOD_HOST_SUPERNODE_BUFFERS] );
/* copy into Lx */
#pragma omp parallel for num_threads(CHOLMOD_OMP_NUM_THREADS) \
private(iidx) if ( nscol2 > 32 )
for ( j=0; j<nscol2; j++ ) {
for ( i=gpu_row_start2*L_ENTRY; i<gpu_row_end*L_ENTRY; i++ ) {
iidx = j*nsrow*L_ENTRY+i;
Lx[psx*L_ENTRY+iidx] = gpu_p->h_Lx[iHostBuff][iidx];
}
}
}
iblock++;
}
/* ---------------------------------------------------------------------- */
/* return */
/* ---------------------------------------------------------------------- */
#ifndef NTIMER
Common->CHOLMOD_GPU_TRSM_TIME += SuiteSparse_time ( ) - tstart ;
#endif
return (1) ;
}
///////////////////////
// Main program entry
///////////////////////
int main(int argc, char** argv)
{
unsigned int max_iters, Nx, Ny, Nz, blockX, blockY, blockZ;
int rank, numberOfProcesses;
if (argc == 8)
{
Nx = atoi(argv[1]);
Ny = atoi(argv[2]);
Nz = atoi(argv[3]);
max_iters = atoi(argv[4]);
blockX = atoi(argv[5]);
blockY = atoi(argv[6]);
blockZ = atoi(argv[7]);
}
else
{
printf("Usage: %s nx ny nz i block_x block_y block_z\n", argv[0]);
exit(1);
}
InitializeMPI(&argc, &argv, &rank, &numberOfProcesses);
AssignDevices(rank);
ECCCheck(rank);
// Define constants
const _DOUBLE_ L = 1.0;
const _DOUBLE_ h = L/(Nx+1);
const _DOUBLE_ dt = h*h/6.0;
const _DOUBLE_ beta = dt/(h*h);
const _DOUBLE_ c0 = beta;
const _DOUBLE_ c1 = (1-6*beta);
// Copy constants to Constant Memory on the GPUs
CopyToConstantMemory(c0, c1);
// Decompose along the z-axis
const int _Nz = Nz/numberOfProcesses;
const int dt_size = sizeof(_DOUBLE_);
// Host memory allocations
_DOUBLE_ *u_new, *u_old;
_DOUBLE_ *h_Uold;
u_new = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(Nz+2));
u_old = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(Nz+2));
if (rank == 0)
{
h_Uold = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(Nz+2));
}
init(u_old, u_new, h, Nx, Ny, Nz);
// Allocate and generate host subdomains
_DOUBLE_ *h_s_Uolds, *h_s_Unews, *h_s_rbuf[numberOfProcesses];
_DOUBLE_ *left_send_buffer, *left_receive_buffer;
_DOUBLE_ *right_send_buffer, *right_receive_buffer;
h_s_Unews = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_Nz+2));
h_s_Uolds = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_Nz+2));
#if defined(DEBUG) || defined(_DEBUG)
if (rank == 0)
{
for (int i = 0; i < numberOfProcesses; i++)
{
h_s_rbuf[i] = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_Nz+2));
checkCuda(cudaHostAlloc((void**)&h_s_rbuf[i], dt_size*(Nx+2)*(Ny+2)*(_Nz+2), cudaHostAllocPortable));
}
}
#endif
right_send_buffer = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_GC_DEPTH));
left_send_buffer = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_GC_DEPTH));
right_receive_buffer = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_GC_DEPTH));
left_receive_buffer = (_DOUBLE_*)malloc(sizeof(_DOUBLE_)*(Nx+2)*(Ny+2)*(_GC_DEPTH));
checkCuda(cudaHostAlloc((void**)&h_s_Unews, dt_size*(Nx+2)*(Ny+2)*(_Nz+2), cudaHostAllocPortable));
checkCuda(cudaHostAlloc((void**)&h_s_Uolds, dt_size*(Nx+2)*(Ny+2)*(_Nz+2), cudaHostAllocPortable));
checkCuda(cudaHostAlloc((void**)&right_send_buffer, dt_size*(Nx+2)*(Ny+2)*(_GC_DEPTH), cudaHostAllocPortable));
checkCuda(cudaHostAlloc((void**)&left_send_buffer, dt_size*(Nx+2)*(Ny+2)*(_GC_DEPTH), cudaHostAllocPortable));
checkCuda(cudaHostAlloc((void**)&right_receive_buffer, dt_size*(Nx+2)*(Ny+2)*(_GC_DEPTH), cudaHostAllocPortable));
checkCuda(cudaHostAlloc((void**)&left_receive_buffer, dt_size*(Nx+2)*(Ny+2)*(_GC_DEPTH), cudaHostAllocPortable));
init_subdomain(h_s_Uolds, u_old, Nx, Ny, _Nz, rank);
// GPU stream operations
cudaStream_t compute_stream;
cudaStream_t data_stream;
checkCuda(cudaStreamCreate(&compute_stream));
checkCuda(cudaStreamCreate(&data_stream));
// GPU Memory Operations
size_t pitch_bytes, pitch_gc_bytes;
_DOUBLE_ *d_s_Unews, *d_s_Uolds;
_DOUBLE_ *d_right_send_buffer, *d_left_send_buffer;
_DOUBLE_ *d_right_receive_buffer, *d_left_receive_buffer;
checkCuda(cudaMallocPitch((void**)&d_s_Uolds, &pitch_bytes, dt_size*(Nx+2), (Ny+2)*(_Nz+2)));
checkCuda(cudaMallocPitch((void**)&d_s_Unews, &pitch_bytes, dt_size*(Nx+2), (Ny+2)*(_Nz+2)));
checkCuda(cudaMallocPitch((void**)&d_left_send_buffer, &pitch_gc_bytes, dt_size*(Nx+2), (Ny+2)*(_GC_DEPTH)));
checkCuda(cudaMallocPitch((void**)&d_left_receive_buffer, &pitch_gc_bytes, dt_size*(Nx+2), (Ny+2)*(_GC_DEPTH)));
checkCuda(cudaMallocPitch((void**)&d_right_send_buffer, &pitch_gc_bytes, dt_size*(Nx+2), (Ny+2)*(_GC_DEPTH)));
checkCuda(cudaMallocPitch((void**)&d_right_receive_buffer, &pitch_gc_bytes, dt_size*(Nx+2), (Ny+2)*(_GC_DEPTH)));
// Copy subdomains from host to device and get walltime
double HtD_timer = 0.;
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
HtD_timer -= MPI_Wtime();
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
checkCuda(cudaMemcpy2D(d_s_Uolds, pitch_bytes, h_s_Uolds, dt_size*(Nx+2), dt_size*(Nx+2), ((Ny+2)*(_Nz+2)), cudaMemcpyDefault));
checkCuda(cudaMemcpy2D(d_s_Unews, pitch_bytes, h_s_Unews, dt_size*(Nx+2), dt_size*(Nx+2), ((Ny+2)*(_Nz+2)), cudaMemcpyDefault));
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
HtD_timer += MPI_Wtime();
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
unsigned int ghost_width = 1;
int pitch = pitch_bytes/dt_size;
int gc_pitch = pitch_gc_bytes/dt_size;
// GPU kernel launch parameters
dim3 threads_per_block(blockX, blockY, blockZ);
unsigned int blocksInX = getBlock(Nx, blockX);
unsigned int blocksInY = getBlock(Ny, blockY);
unsigned int blocksInZ = getBlock(_Nz-2, k_loop);
dim3 thread_blocks(blocksInX, blocksInY, blocksInZ);
dim3 thread_blocks_halo(blocksInX, blocksInY);
//MPI_Status status;
MPI_Status status[numberOfProcesses];
MPI_Request gather_send_request[numberOfProcesses];
MPI_Request right_send_request[numberOfProcesses], left_send_request[numberOfProcesses];
MPI_Request right_receive_request[numberOfProcesses], left_receive_request[numberOfProcesses];
double compute_timer = 0.;
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
compute_timer -= MPI_Wtime();
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
for(unsigned int iterations = 0; iterations < max_iters; iterations++)
{
// Compute right boundary data on device 0
if (rank == 0) {
int kstart = (_Nz+1)-ghost_width;
int kstop = _Nz+1;
ComputeInnerPointsAsync(thread_blocks_halo, threads_per_block, data_stream, d_s_Unews, d_s_Uolds, pitch, Nx, Ny, _Nz, kstart, kstop);
CopyBoundaryRegionToGhostCellAsync(thread_blocks_halo, threads_per_block, data_stream, d_s_Unews, d_right_send_buffer, Nx, Ny, _Nz, pitch, gc_pitch, 0);
checkCuda(cudaMemcpy2DAsync(right_send_buffer, dt_size*(Nx+2), d_right_send_buffer, pitch_gc_bytes, dt_size*(Nx+2), (Ny+2)*(_GC_DEPTH), cudaMemcpyDefault, data_stream));
checkCuda(cudaStreamSynchronize(data_stream));
MPI_CHECK(MPI_Isend(right_send_buffer, (Nx+2)*(Ny+2)*(_GC_DEPTH), MPI_DOUBLE, 1, 0, MPI_COMM_WORLD, &right_send_request[rank]));
}
else
{
int kstart = 1;
int kstop = 1+ghost_width;
ComputeInnerPointsAsync(thread_blocks_halo, threads_per_block, data_stream, d_s_Unews, d_s_Uolds, pitch, Nx, Ny, _Nz, kstart, kstop);
CopyBoundaryRegionToGhostCellAsync(thread_blocks_halo, threads_per_block, data_stream, d_s_Unews, d_left_send_buffer, Nx, Ny, _Nz, pitch, gc_pitch, 1);
checkCuda(cudaMemcpy2DAsync(left_send_buffer, dt_size*(Nx+2), d_left_send_buffer, pitch_gc_bytes, dt_size*(Nx+2), (Ny+2)*(_GC_DEPTH), cudaMemcpyDefault, data_stream));
checkCuda(cudaStreamSynchronize(data_stream));
MPI_CHECK(MPI_Isend(left_send_buffer, (Nx+2)*(Ny+2)*(_GC_DEPTH), MPI_DOUBLE, 0, 1, MPI_COMM_WORLD, &left_send_request[rank]));
}
// Compute inner nodes for device 0
if (rank == 0) {
int kstart = 1;
int kstop = (_Nz+1)-ghost_width;
ComputeInnerPointsAsync(thread_blocks, threads_per_block, compute_stream, d_s_Unews, d_s_Uolds, pitch, Nx, Ny, _Nz, kstart, kstop);
}
// Compute inner nodes for device 1
else
{
int kstart = 1+ghost_width;
int kstop = _Nz+1;
ComputeInnerPointsAsync(thread_blocks, threads_per_block, compute_stream, d_s_Unews, d_s_Uolds, pitch, Nx, Ny, _Nz, kstart, kstop);
}
// Receive data from device 1
if (rank == 0) {
MPI_CHECK(MPI_Irecv(left_send_buffer, (Nx+2)*(Ny+2)*(_GC_DEPTH), MPI_DOUBLE, 1, 1, MPI_COMM_WORLD, &right_receive_request[rank]));
}
else
{
MPI_CHECK(MPI_Irecv(right_send_buffer, (Nx+2)*(Ny+2)*(_GC_DEPTH), MPI_DOUBLE, 0, 0, MPI_COMM_WORLD, &left_receive_request[rank]));
}
if (rank == 0) {
MPI_CHECK(MPI_Wait(&right_receive_request[rank], &status[rank]));
checkCuda(cudaMemcpy2DAsync(d_right_receive_buffer, pitch_gc_bytes, left_send_buffer, dt_size*(Nx+2), dt_size*(Nx+2), ((Ny+2)*(_GC_DEPTH)), cudaMemcpyDefault, data_stream));
CopyGhostCellToBoundaryRegionAsync(thread_blocks_halo, threads_per_block, data_stream, d_s_Unews, d_right_receive_buffer, Nx, Ny, _Nz, pitch, gc_pitch, 0);
}
else
{
MPI_CHECK(MPI_Wait(&left_receive_request[rank], &status[rank]));
checkCuda(cudaMemcpy2DAsync(d_left_receive_buffer, pitch_gc_bytes, right_send_buffer, dt_size*(Nx+2), dt_size*(Nx+2), ((Ny+2)*(_GC_DEPTH)), cudaMemcpyDefault, data_stream));
CopyGhostCellToBoundaryRegionAsync(thread_blocks_halo, threads_per_block, data_stream, d_s_Unews, d_left_receive_buffer, Nx, Ny, _Nz, pitch, gc_pitch, 1);
}
if (rank == 0)
{
MPI_CHECK(MPI_Wait(&right_send_request[rank], MPI_STATUS_IGNORE));
}
else
{
MPI_CHECK(MPI_Wait(&left_send_request[rank], MPI_STATUS_IGNORE));
}
// Swap pointers on the host
checkCuda(cudaDeviceSynchronize());
swap(_DOUBLE_*, d_s_Unews, d_s_Uolds);
}
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
compute_timer += MPI_Wtime();
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
// Copy data from device to host
double DtH_timer = 0;
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
DtH_timer -= MPI_Wtime();
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
checkCuda(cudaMemcpy2D(h_s_Uolds, dt_size*(Nx+2), d_s_Uolds, pitch_bytes, dt_size*(Nx+2), (Ny+2)*(_Nz+2), cudaMemcpyDefault));
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
DtH_timer += MPI_Wtime();
MPI_CHECK(MPI_Barrier(MPI_COMM_WORLD));
// Gather results from subdomains
MPI_CHECK(MPI_Isend(h_s_Uolds, (Nx+2)*(Ny+2)*(_Nz+2), MPI_DOUBLE, 0, 0, MPI_COMM_WORLD, &gather_send_request[rank]));
if (rank == 0)
{
for (int i = 0; i < numberOfProcesses; i++)
{
MPI_CHECK(MPI_Recv(h_s_rbuf[i], (Nx+2)*(Ny+2)*(_Nz+2), MPI_DOUBLE, i, 0, MPI_COMM_WORLD, &status[rank]));
merge_domains(h_s_rbuf[i], h_Uold, Nx, Ny, _Nz, i);
}
}
// Calculate on host
#if defined(DEBUG) || defined(_DEBUG)
if (rank == 0)
{
cpu_heat3D(u_new, u_old, c0, c1, max_iters, Nx, Ny, Nz);
}
#endif
if (rank == 0)
{
float gflops = CalcGflops(compute_timer, max_iters, Nx, Ny, Nz);
PrintSummary("3D Heat (7-pt)", "Plane sweeping", compute_timer, HtD_timer, DtH_timer, gflops, max_iters, Nx);
_DOUBLE_ t = max_iters * dt;
CalcError(h_Uold, u_old, t, h, Nx, Ny, Nz);
}
Finalize();
// Free device memory
checkCuda(cudaFree(d_s_Unews));
checkCuda(cudaFree(d_s_Uolds));
checkCuda(cudaFree(d_right_send_buffer));
checkCuda(cudaFree(d_left_send_buffer));
checkCuda(cudaFree(d_right_receive_buffer));
checkCuda(cudaFree(d_left_receive_buffer));
// Free host memory
checkCuda(cudaFreeHost(h_s_Unews));
checkCuda(cudaFreeHost(h_s_Uolds));
#if defined(DEBUG) || defined(_DEBUG)
if (rank == 0)
{
for (int i = 0; i < numberOfProcesses; i++)
{
checkCuda(cudaFreeHost(h_s_rbuf[i]));
}
free(h_Uold);
}
#endif
checkCuda(cudaFreeHost(left_send_buffer));
checkCuda(cudaFreeHost(left_receive_buffer));
checkCuda(cudaFreeHost(right_send_buffer));
checkCuda(cudaFreeHost(right_receive_buffer));
checkCuda(cudaDeviceReset());
free(u_old);
free(u_new);
return 0;
}
cudaError_t WINAPI wine_cudaMemcpy2DAsync( void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream ) {
WINE_TRACE("\n");
return cudaMemcpy2DAsync( dst, dpitch, src, spitch, width, height, kind, stream );
}
void magmablas_chemm_mgpu_com(
char side, char uplo, magma_int_t m, magma_int_t n,
magmaFloatComplex alpha,
magmaFloatComplex *dA[], magma_int_t ldda, magma_int_t offset,
magmaFloatComplex *dB[], magma_int_t lddb,
magmaFloatComplex beta, magmaFloatComplex *dC[], magma_int_t lddc,
magmaFloatComplex *dwork[], magma_int_t dworksiz,
magmaFloatComplex *C, magma_int_t ldc,
magmaFloatComplex *work[], magma_int_t worksiz,
magma_int_t ngpu, magma_int_t nb,
magma_queue_t streams[][20], magma_int_t nstream,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10], magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
//printf("####################################################\n");
//printf(" start chemm \n");
//printf("####################################################\n");
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nstream >= ngpu );
assert( nbevents >= ngpu*ngpu );
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *dwork1[MagmaMaxGPUs];
magmaFloatComplex *dwork2[MagmaMaxGPUs];
magma_int_t maxgsize = n*m;
magma_int_t lddwork = lddc;
magma_int_t ldwork = m;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev]; // size of dwork1 is n*lddwork
dwork2[dev] = dwork[dev]+n*lddwork; // size of dwork2 is maxgsize*ngpu
}
assert( dworksiz >= (n*lddwork+maxgsize*ngpu) );
assert( worksiz >= (n*ldwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev, devperm, myblk, mycolsize, myblkoffst;
magma_int_t gmaster;
magma_int_t masterdev, lcdev, lccolsize, myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m, nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
cudaMemset(dwork(dev,0,0), 0, (lddwork)*(n)*sizeof(magmaFloatComplex) );
// put all dC on all dev to 0 except the one which
// hold i==0 because this one has to multiply by beta.
if(dev!=stdev){
cudaMemset(dC(dev,0,0), 0, (lddc)*(n)*sizeof(magmaFloatComplex) );
}
}
magma_int_t newoffset = offset;
// 1. symmetrize
if(blockoffset>0){
newoffset = offset+fstblksiz; // newoffset is adjusted over nb
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > stdev?1:0);
//printf("STDEV %d voici offset %d remm %d myblockoffset %d siz %d \n", stdev, offset, remm, myblkoffst, fstblksiz);
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
magmablas_csymmetrize_tiles( MagmaLower, fstblksiz, dA(stdev, offset, myblkoffst*nb+blockoffset), ldda, 1, ngpu*nb, nb );
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t newstdev = (newoffset/nb)%ngpu;
magma_int_t nbblk = remm/nb; // number of block of size nb. if m%nb>0 then a last block exist and is of size ib=m%nb
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-newstdev+ngpu)%ngpu) ? 1:0 );
magma_int_t devperm = (dev-newstdev+ngpu)%ngpu;
magma_int_t nbblkoffst = newoffset/nb;
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
//printf("dev %d devperm %d newoffset %d rowoff %d coloff %d myblk %d \n", dev, devperm, newoffset, newoffset+devperm*nb, myblkoffst*nb, myblk);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magmablas_csymmetrize_tiles( MagmaLower, nb, dA(dev, newoffset+devperm*nb, myblkoffst*nb), ldda, myblk, ngpu*nb, nb );
if(remm%nb>0){
magma_int_t nblstblks = (nbblk+1)%ngpu;
magma_int_t devlstblk = (nblstblks-1+ngpu)%ngpu;
//printf("==> siz %d devperm %d, devlstblk %d, newoffset+nbblk*nb %d, myblkoffst*nb+ myblk*nb %d\n", remm % nb, devperm, devlstblk, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb);
if(devperm==devlstblk)
magmablas_csymmetrize( MagmaLower, remm % nb, dA(dev, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb), ldda ); // last partial tile
}
}
/*
magma_int_t siz = m+offset;
magmaFloatComplex *R=(magmaFloatComplex *) malloc(siz*siz*sizeof(magmaFloatComplex));
// collecte back A
magmablas_cgetmatrix_1D_bcyclic( siz, siz, dA, ldda, R, siz, ngpu, nb );
magma_setdevice( 0 );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
//magma_cgetmatrix( siz, siz, dA[0], ldda, R, siz );
FILE *trace_file;
trace_file = fopen("AJETE/Aafter", "w");
for (int j = 0; j < siz ; j++)
for (int i = 0; i < siz ; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, R[j*siz+i]);
fclose(trace_file);
return;
*/
// ROW GEMM transpose a row and make a gemm with a block
// if only 1 GPU used the ROW GEMM is integrated with the
// COL GEMM (better accuracy observed) and better perf
if(ngpu>1){
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
//magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nbblk = magma_ceildiv(i, nb);
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-stdev+ngpu)%ngpu) ? 1:0 );
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
magma_int_t myrowsize = myblk * nb;
magma_int_t coloffset = myblkoffst*nb;
if(dev==stdev) {
myrowsize = myrowsize -blockoffset;
coloffset = myblkoffst*nb+blockoffset;
}
//printf("ROW GEMM: voici i %d ib %d ioff %d nbblkoffst %d stdev %d dev %d myblk %d myblkoffset %d coloffset %d rowsize %d\n", i, ib, ioff, nbblkoffst, stdev, dev, myblk, myblkoffst, coloffset, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_cgemm( MagmaConjTrans, MagmaNoTrans, myrowsize, n, ib,
alpha, dA(dev,ioff,coloffset), ldda,
dB(dev,i,0), lddb,
c_one, dwork(dev,0,0), lddwork );
}
}
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_event_record(redevents[dev][1], streams[dev][1]);
}
}
// COL GEMM
// blockoffset is offset within first block; for subsequent blocks it is 0
if(blockoffset>0){
magma_int_t ib = min( nb-blockoffset, m ); // block size
magma_int_t iblock = (offset / nb) / ngpu; // local block id
magma_int_t di = iblock*nb+blockoffset; // local index in parent matrix
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
//printf("DEV %d COL GEMM first ioff %d di %d m %d n %d ib %d \n", stdev, offset, di, m, n, ib);
magma_cgemm( MagmaNoTrans, MagmaNoTrans, m, n, ib,
alpha, dA(stdev,offset,di), ldda,
dB(stdev,0,0), lddb,
beta, dC(stdev,0,0), lddc );
}
// COL GEMM
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
magma_int_t iblock = (ioff / nb) / ngpu; // local block id
magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t di = iblock*nb; // local index in parent matrix
//printf("DEV %d COL GEMM i %d ioff %d di %d m-i %d n %d ib %d \n", dev, i, ioff, di, m-i, n, ib);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
if(i==0){
magma_cgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
beta, dC(dev,i,0), lddc );
}else{
magma_cgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,i,0), lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
// if only 1 GPU is used, do the ROW GEMM
if(ngpu==1){
// NOTE THAT because the COL gemm write dC below the diagonal (i)
// and the ROW GEMM write dC from 0 to diag-1, so they could
// run in parallel on diferent stream.
//
// NO NO NO because
// it might happen that col finished i and strated i+1 while row still at i
// magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_cgemm( MagmaConjTrans, MagmaNoTrans, i, n, ib,
alpha, dA(dev,ioff,offset), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,0,0), lddc );
}
}
if(ngpu>1){
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t nbblkrow = nbblk-1;
magma_int_t devperm = (dev-stdev+ngpu)%ngpu;
magma_int_t myblk = (nbblkrow/ngpu) + (nbblkrow%ngpu > devperm ? 1:0 );
magma_int_t myrowsize = myblk * nb;
if(dev==stdev) {
myrowsize = myrowsize - blockoffset;
}
//printf("blockoffset %d nbblkrow %d devperm %d DEV %d RECEIVING myblk %d myrowsize %d\n", blockoffset, nbblkrow, devperm, dev, myblk, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][1]);
//magma_queue_sync( streams[ dev ][ 1 ] );
// for each dev add the computed ROW block each on its placment with dC
for( magma_int_t blki = 0; blki < myblk; ++blki){
magma_int_t gbblki = (blki*ngpu + devperm)*nb - blockoffset;
magma_int_t lcblki = blki*nb;
magma_int_t ib = nb;// min(nb, m-gbblki);
if(dev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
magmablas_cgeadd(ib, n, c_one,
&dwork[dev][lcblki], lddwork,
&dC[dev][gbblki] , lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
}
}
}
// ===========================================================
// COMMUNICATION ALL_REDUCE_SUM
// ===========================================================
if(ngpu==1){
return;
}
// INITIALIZE COMM
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize, m);
if(mycolsize>0){
gpuisactive[dev] = mycolsize;
if(masterdev==-1) {
masterdev = dev;
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
cudaDeviceSynchronize();
}
*/
//*******************************
// each GPU send its result
// to its master. The master make
// the addition and then send to
// to the masters of other complex
// and receive from the masters of
// other complex make the addition
// and broadcast locally the final
// result.
//*******************************
//printf("=======================================================================\n");
//printf(" sending to my master \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if complex is active
if(masterdev!=-1){
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU. if I am not the master, then send my result to my master.
// store result on dwork[masterdev][dev*maxgsize]
if(dev!=masterdev){
magma_setdevice( dev );
//printf(" GPU %d sending to my master %d\n", dev, masterdev);
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][0]);
// sending to the master of my complex
cudaMemcpy2DAsync(&dwork2[masterdev][maxgsize*dev], m*sizeof(magmaFloatComplex),
&dC[dev][0], lddc*sizeof(magmaFloatComplex),
m*sizeof(magmaFloatComplex), n,
cudaMemcpyDeviceToDevice, streams[dev][0]);
magma_event_record(redevents[dev][masterdev], streams[dev][0]);
} // end I am not the masterdev
}// end if mycolsize>0
}// for idev
}// end of if masterdev!=-1 maening complex is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
cudaDeviceSynchronize();
}
*/
//printf("=======================================================================\n");
//printf(" each master do addition of local result and broadcast to other masters \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if complex is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][0]);
// ========================================
// local addition
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" master %d receiving from %d and adding \n", masterdev, lcdev);
// this is an active GPU of my complex.
// wait I received what he send it to me and then do addition.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[lcdev][masterdev]);
magmablas_cgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*lcdev], m,
&dC[masterdev][0] , lddc );
}
}// for l=1:myngpu
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
//
// ========================================
// send to other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //complex is active
//Master has to wait until finish the local addition then send using gmaster stream.
//use stream 0 to make it sequential or stream gmaster to make it parallel.
//Now both re the same.
//printf(" master %d from cmplx %d sending to other master %d on cmplx %d \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ gmaster ], redevents[masterdev][masterdev]);
cudaMemcpy2DAsync(&dwork2[gmaster][maxgsize*masterdev], m*sizeof(magmaFloatComplex),
&dC[masterdev][0], lddc*sizeof(magmaFloatComplex),
m*sizeof(magmaFloatComplex), n,
cudaMemcpyDeviceToDevice, streams[masterdev][gmaster]);
magma_event_record(redevents[masterdev][gmaster], streams[masterdev][gmaster]);
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][gmaster]);
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// ========================================
}// end of if masterdev!=-1 maening complex is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
cudaDeviceSynchronize();
}
*/
//printf("=======================================================================\n");
//printf(" each master wait receiving other masters results, do the addition and broadcast locally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if complex is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// master has to wait until finishing all the send to other masters.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
// ========================================
// addition of results from other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //complex is active
//Master has to wait until receiving from gmaster, then do addition using stream 0
//printf(" master %d from cmplx %d receiving from other master %d on cmplx %d and adding \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[gmaster][masterdev]);
magmablas_cgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*gmaster], m,
&dC[masterdev][0] , lddc );
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
// ========================================
// local broadcast of final results
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
// this is an active GPU of my complex.
// wait the previous addition is done maening stream 0 is finished and broadcast sequentially for now.
// to make it parallel put stream lcdev instead of stream 0
//printf(" master %d broadcasting local to %d \n", masterdev, lcdev);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
cudaMemcpy2DAsync(&dC[lcdev][0], lddc*sizeof(magmaFloatComplex),
&dC[masterdev][0], lddc*sizeof(magmaFloatComplex),
m*sizeof(magmaFloatComplex), n,
cudaMemcpyDeviceToDevice, streams[masterdev][0]);
magma_event_record(redevents[masterdev][lcdev], streams[masterdev][0]);
}
}// for l=1:myngpu
// ========================================
}// end of if masterdev!=-1 maening complex is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
cudaDeviceSynchronize();
}
*/
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if complex is active
if(masterdev!=-1){
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if(lccolsize>0){
magma_setdevice( lcdev );
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[lcdev][0]);
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[masterdev][lcdev]);
}
}// for l=1:myngpu
}// end of if masterdev!=-1 maening complex is active
}// for cmplxid
//printf("****************************************************\n");
//printf(" finish chemm \n");
//printf("****************************************************\n");
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}