/*
Usage:
cutorch.streamWaitFor(waiterStream, {waitForStream1, ..., waitForStreamN})
for streams on the current device. Creates a one-way barrier where
waiterStream waits for waitForStream1-N to reach the current point.
*/
static int cutorch_streamWaitFor(lua_State *L)
{
THCState *state = cutorch_getstate(L);
int curDev = -1;
THCudaCheck(cudaGetDevice(&curDev));
/* Check that the waiting stream is in bounds; this will error out if not */
int waitingId = (int) luaL_checknumber(L, 1);
cudaStream_t streamWaiting =
THCState_getDeviceStream(state, curDev, waitingId);
/* Validate the streams that we are waiting on */
int streams = checkAndCountListOfStreams(L, state, 2, curDev);
if (streams < 1) {
/* nothing to synchronize */
return 0;
}
/* One-way dependency; streamWaiting will wait for the list of streams to
wait on to complete execution of pending scheduled kernels/events */
cudaEvent_t * events = (cudaEvent_t*)malloc(sizeof(cudaEvent_t) * streams);
createSingleDeviceEvents(L, state, 2, curDev, events);
/* Then, wait on them */
for (int i = 0; i < streams; i++) {
THCudaCheck(cudaStreamWaitEvent(streamWaiting, events[i], 0));
THCudaCheck(cudaEventDestroy(events[i]));
}
free(events);
return 0;
}
static int cutorch_Event_waitOn(lua_State *L)
{
cudaEvent_t *event = luaT_checkudata(L, 1, "cutorch.Event");
THCState *state = cutorch_getstate(L);
THCudaCheck(cudaStreamWaitEvent(THCState_getCurrentStream(state), *event, 0));
return 0;
}
/*
Usage:
cutorch.streamWaitForMultiDevice(gpuWaiter, streamWaiter,
{[gpu1]={stream1_1, ..., stream1_N},
[gpuK]={streamK_1, ..., streamK_M}})
with a specified GPU per each list of streams.
Stream (gpuWaiter, streamWaiter) will wait on all of the other streams
(gpu1, stream1_1), ..., (gpu1, stream1_N), ...,
(gpuK, streamK_1), ..., (gpuK, streamK_M) to complete fully, as a one-way
barrier only (only streamWaiter is blocked).
The streams to wait on are bucketed per device. Equivalent to
streamWaitFor() if only one GPU's streams are listed.
*/
static int cutorch_streamWaitForMultiDevice(lua_State *L)
{
THCState *state = cutorch_getstate(L);
int prevDev = -1;
THCudaCheck(cudaGetDevice(&prevDev));
/* Validate waiting (gpu, stream); this will error out if not */
int gpuWaiter = (int) luaL_checknumber(L, 1) - 1;
int streamWaiter = (int) luaL_checknumber(L, 2);
cudaStream_t streamWaiting =
THCState_getDeviceStream(state, gpuWaiter, streamWaiter);
/* Validate and count set of {gpu={streams...}} we are waiting on */
int gpus = 0;
int streams = 0;
checkAndCountListOfGPUStreamPairs(L, state, 3, &gpus, &streams);
if (streams < 1) {
/* nothing to synchronize together */
return 0;
}
/*
Events can only be recorded on the same device on which they are created.
-For each GPU, create and record event per each stream given
for that GPU.
-For (gpuWaiter, streamWaiter), wait on all of the above events.
*/
cudaEvent_t* events = (cudaEvent_t*) malloc(sizeof(cudaEvent_t) * streams);
/* First, create an event per GPU and record events for the specified stream
on that GPU */
createMultiDeviceEvents(L, state, 3, events);
/* Then, wait on the events */
THCudaCheck(cudaSetDevice(gpuWaiter));
for (int i = 0; i < streams; ++i) {
THCudaCheck(cudaStreamWaitEvent(streamWaiting, events[i], 0));
}
/* Clean up events */
for (int i = 0; i < streams; ++i) {
THCudaCheck(cudaEventDestroy(events[i]));
}
free(events);
THCudaCheck(cudaSetDevice(prevDev));
return 0;
}
void waitSingleDeviceEvents(lua_State *L, THCState *state, int arg,
int device, cudaEvent_t * event, int numEvents)
{
/* Push table to top */
lua_pushvalue(L, arg);
/* Then, wait on the events. Each stream is actually waiting on itself here
too, but that's harmless and isn't worth weeding out. */
lua_pushnil(L);
while (lua_next(L, -2)) {
int streamId = (int) lua_tonumber(L, -1);
cudaStream_t stream =
THCState_getDeviceStream(state, device, streamId);
for (int i = 0; i < numEvents; i++) {
THCudaCheck(cudaStreamWaitEvent(stream, event[i], 0));
}
lua_pop(L, 1);
}
/* Pop table from top */
lua_pop(L, 1);
}
__host__ __device__
void checked_launch_kernel_after_event(void* kernel, ::dim3 grid_dim, ::dim3 block_dim, int shared_memory_size, cudaStream_t stream, cudaEvent_t event, const Args&... args)
{
#if __cuda_lib_has_cudart
if(event)
{
// make the next launch wait on the event
throw_on_error(cudaStreamWaitEvent(stream, event, 0), "cuda::detail::checked_launch_kernel_after_event(): cudaStreamWaitEvent()");
}
#else
// the error message we return depends on how the program was compiled
const char* error_message =
# ifndef __CUDA_ARCH__
"cuda::detail::checked_launch_kernel_after_event(): CUDA kernel launch from host requires nvcc"
# else
"cuda::detail::checked_launch_kernel_after_event(): CUDA kernel launch from device requires arch=sm_35 or better and rdc=true"
# endif
;
throw_on_error(cudaErrorNotSupported, error_message);
#endif // __cuda_lib_has_cudart
checked_launch_kernel(kernel, grid_dim, block_dim, shared_memory_size, stream, args...);
}
void waitMultiDeviceEvents(lua_State *L, THCState *state, int arg,
cudaEvent_t* events, int streams)
{
/* Push {gpu={streams...}} table */
lua_pushvalue(L, arg);
/* Then, wait on the events. Each stream is actually waiting on itself here
too, but that's harmless and isn't worth weeding out. */
lua_pushnil(L);
while (lua_next(L, -2)) {
int device = (int) lua_tonumber(L, -2) - 1;
THCudaCheck(cudaSetDevice(device));
/* Push stream table */
lua_pushvalue(L, -1);
lua_pushnil(L);
while (lua_next(L, -2)) {
int streamId = (int) lua_tonumber(L, -1);
cudaStream_t stream =
THCState_getDeviceStream(state, device, streamId);
/* Each stream waits on all events */
for (int i = 0; i < streams; ++i) {
THCudaCheck(cudaStreamWaitEvent(stream, events[i], 0));
}
lua_pop(L, 1);
}
/* Pop stream table and GPU entry */
lua_pop(L, 2);
}
/* Pop {gpu={streams...}} table */
lua_pop(L, 1);
}
void GranularGPUDataTransferer::WaitForSyncPointOnAssignStreamAsync()
{
PrepareDevice(m_deviceId);
cudaStreamWaitEvent(m_assignStream, m_syncEvent, 0 /*flags 'must be 0'*/) || "cudaStreamWaitEvent failed";
}
extern "C" magma_int_t
magma_sbulge_applyQ_v2_m(magma_int_t ngpu, magma_side_t side,
magma_int_t NE, magma_int_t N,
magma_int_t NB, magma_int_t Vblksiz,
float *E, magma_int_t lde,
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 MAGMA_SUCCESS;
}
if ( N == 0 ) {
return MAGMA_SUCCESS;
}
if ( NB == 0 ) {
return MAGMA_SUCCESS;
}
/* ==========================================
* 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
cudaDeviceSynchronize();
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t nbevents =2, nstream=2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t myevent[MagmaMaxGPUs][20];
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t i = 0; i < nstream; ++i ) {
cudaStreamCreate( &streams[dev][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&myevent[dev][i],cudaEventDisableTiming);
}
}
// Azzam 21/11/2012
// NOTE THAT dwork was of size 2*NE*Vblksiz+...
// but I am thinking why not modifing it to NE*Vblksiz+...
// BUT NO because the 2* is used because of making 2 streams working and so
// they might be using dwork in parallel
float *dE[MagmaMaxGPUs];
float *dwork[MagmaMaxGPUs], *dwork0[MagmaMaxGPUs], *dwork1[MagmaMaxGPUs];
//float *dwvt[MagmaMaxGPUs];
float *dwvt0[MagmaMaxGPUs], *dwvt1[MagmaMaxGPUs];
float *dT0[MagmaMaxGPUs], *dV0[MagmaMaxGPUs], *dT1[MagmaMaxGPUs], *dV1[MagmaMaxGPUs];
magma_int_t dev;
magma_int_t ldde = N;
magma_int_t lddv = ldv;
magma_int_t lddt = ldt;
magma_int_t ne_loc = magma_ceildiv(NE, ngpu);
if (ne_loc < 256)
ne_loc=256;
magma_int_t dwVTsiz = lddv*Vblksiz; // lddv*lddv + lddv*NE; // lddv*Vblksiz;
magma_int_t dworksiz = ne_loc*Vblksiz; // lddv*Vblksiz; // NE*Vblksiz;
ngpu = min(ngpu, magma_ceildiv(NE,ne_loc)); // Don't use GPU that will not have data.
// copy dE to GPUs
for (dev=0; dev < ngpu; ++dev) {
magma_setdevice( dev );
if (MAGMA_SUCCESS != magma_smalloc( &dE[dev], ldde * ne_loc)) {
printf ("!!!! magma_sbulge_applyQ magma_alloc failed for: dE\n" );
exit(-1);
}
if (MAGMA_SUCCESS != magma_smalloc( &dwork[dev], 2*dworksiz + 2*dwVTsiz + 2*Vchunksiz* (Vblksiz* (lddv+lddt)) )) {
printf ("!!!! magma_sbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
dwork0[dev] = dwork[dev]; // size = dworksiz;
dwork1[dev] = dwork0[dev] + dworksiz; // size = dworksiz;
dwvt0[dev] = dwork[dev] + 2*dworksiz; // size = dwVTsiz;
dwvt1[dev] = dwvt0[dev] + dwVTsiz; // size = dwVTsiz;
dV0[dev] = dwork[dev] + 2*dworksiz + 2*dwVTsiz;
dT0[dev] = dV0[dev] + Vchunksiz*Vblksiz*lddv;
dV1[dev] = dT0[dev] + Vchunksiz*Vblksiz*lddt;
dT1[dev] = dV1[dev] + Vchunksiz*Vblksiz*lddv;
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_ssetmatrix_async( N, ie_loc, E+lde*ne_loc*dev, lde, dE(dev, 0, 0), ldde, streams[dev][1] );
}
// make overlapped copy
magma_int_t ncpy = 0;
magma_int_t copyed=0, copyst=0;
magma_int_t blkcnt,nothing, mysiz, flip, vld,tld, locpos;
findVTsiz(N, NB, Vblksiz, &blkcnt, ¬hing);
flip = 0;
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
#ifdef ENABLE_DEBUG
printf(" APPLY Q_v22_m GPU with NGPU %d N %d, NE %d, NB %d, Vblksiz %d, versionL %d versionR %d SIDE %c \n",
ngpu, N, NE, NB, Vblksiz, versionL, versionR, side);
#endif
/*
* MagmamaLeft
*/
if (side == MagmaLeft) {
/*
* Version 113:
* loop over the block_col (nt) and for each find the
* number of tiles (mt) in this block_col. then loop over mt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
if ( versionL == 113 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=nt-1; blkj >= 0; blkj--) {
/* the index of the first row on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=mt; blki > 0; blki--) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d vpos %d \n",blkj,blki,Vm, Vn,mycol,vpos);
magma_bulge_findpos113(N, NB, Vblksiz, mycol, myrow, &blkid);
// COPY Vchunksiz Vs and Vchunksiz Ts to GPU and store it in dV0/dV1 and dT0/dT1
if (ncpy == 0) {
// flip = 1 for this.
copyst = 0; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
if (mysiz > 0) {
ncpy = 1;
flip = 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, streams[dev][1]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, streams[dev][1]);
}
//printf("doing the first copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
}
}
if (blkid == copyst) {
flip = ncpy % 2;
copyst = copyed; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
//printf(" get to copy blkid %d blkid+(2*Vchunksiz) %d copyst %d copyed %d\n",blkid,blkid+(Vchunksiz),copyst,copyed);
if (mysiz > 0) {
ncpy = ncpy + 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
if (flip == 0) { // now I am working on dV0 so copy the next and put it on dV1
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, streams[dev][1]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, streams[dev][1]);
}
} else { // now I am working on dV1 so copy the next and put it on dV0
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV0 dT0\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV0[dev], vld, streams[dev][0]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT0[dev], tld, streams[dev][0]);
}
}
}
}
if ((Vm > 0) && (Vn > 0)) {
locpos = blkid%Vchunksiz;
magma_int_t lcvpos = locpos*Vblksiz*lddv;
magma_int_t lctpos = locpos*Vblksiz*lddt;
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d locvpos %5d loctpos %5d blkid %2d using data in dV%1d dT%1d \n",blkj,blki,Vm, Vn,mycol,lcvpos,lctpos, blkid,flip,flip);
if (flip == 0) {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
cudaStreamWaitEvent(streams[dev][0], myevent[dev][1], 0);
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib);
magma_slarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib, dwvt0[dev], Vm);
}
cudaEventRecord(myevent[dev][0], streams[dev][0]);
}
} else {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][1]);
cudaStreamWaitEvent(streams[dev][1], myevent[dev][0], 0);
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib);
magma_slarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib, dwvt1[dev], Vm);
}
cudaEventRecord(myevent[dev][1], streams[dev][1]);
}
}
} // end for (Vm &Vn) > 0
} // end for blki
} // end for blkj
} // end if version=113
/*
* Version 114:
* loop over the block_row (mt) and for each find diagonally the
* number of tiles (nt) in this block_row. then loop over nt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
else {
printf("versionL 114 not implemented in sbulge_applyQ_v2_m\n");
exit(-1);
mt = magma_ceildiv((N-1),NB);
for (blki = mt; blki > 0; blki--) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = nt-1; blkj >= 0; blkj--) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_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 {
printf("Side 'R' not implemented in sbulge_applyQ_v2_m\n");
exit(-1);
/*
* Version 91:
*/
if ( versionR == 91 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=0; blkj < nt; blkj++) {
/* the index of the first myrow on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=1; blki <= mt; blki++) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( (blkj == nt-1) && (blki == mt) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
mycol = blkj*Vblksiz;
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_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
// copy back the dE form each GPU
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
cudaStreamWaitEvent(streams[dev][0], myevent[dev][1], 0);
cudaStreamWaitEvent(streams[dev][0], myevent[dev][0], 0);
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_sgetmatrix_async( N, ie_loc, dE(dev, 0, 0), ldde, E+lde*ne_loc*dev, lde, streams[dev][0] );
cudaEventRecord(myevent[dev][0], streams[dev][0]);
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
cudaStreamWaitEvent(streams[dev][0], myevent[dev][0], 0);
cudaDeviceSynchronize(); // no need for synchronize
magma_free(dwork[dev]);
magma_free(dE[dev]);
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventDestroy(myevent[dev][i]);
}
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[dev][i] );
}
}
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
return MAGMA_SUCCESS;
}
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_updateC))
(
Int ndrow1, /* C is ndrow2-by-ndrow2 */
Int ndrow2,
Int ndrow, /* leading dimension of Lx */
Int ndcol, /* L1 is ndrow1-by-ndcol */
Int nsrow,
Int pdx1, /* L1 starts at Lx + L_ENTRY*pdx1 */
/* L2 starts at Lx + L_ENTRY*(pdx1 + ndrow1) */
Int pdi1,
double *Lx,
double *C,
cholmod_common *Common,
cholmod_gpu_pointers *gpu_p
)
{
double *devPtrLx, *devPtrC ;
double alpha, beta ;
cublasStatus_t cublasStatus ;
cudaError_t cudaStat [2] ;
Int ndrow3 ;
int icol, irow;
int iHostBuff, iDevBuff ;
#ifndef NTIMER
double tstart = 0;
#endif
if ((ndrow2*L_ENTRY < CHOLMOD_ND_ROW_LIMIT) ||
(ndcol*L_ENTRY < CHOLMOD_ND_COL_LIMIT))
{
/* too small for the CUDA BLAS; use the CPU instead */
return (0) ;
}
ndrow3 = ndrow2 - ndrow1 ;
#ifndef NTIMER
Common->syrkStart = SuiteSparse_time ( ) ;
Common->CHOLMOD_GPU_SYRK_CALLS++ ;
#endif
/* ---------------------------------------------------------------------- */
/* allocate workspace on the GPU */
/* ---------------------------------------------------------------------- */
iHostBuff = (Common->ibuffer)%CHOLMOD_HOST_SUPERNODE_BUFFERS;
iDevBuff = (Common->ibuffer)%CHOLMOD_DEVICE_STREAMS;
/* cycle the device Lx buffer, d_Lx, through CHOLMOD_DEVICE_STREAMS,
usually 2, so we can overlap the copy of this descendent supernode
with the compute of the previous descendant supernode */
devPtrLx = (double *)(gpu_p->d_Lx[iDevBuff]);
/* very little overlap between kernels for difference descendant supernodes
(since we enforce the supernodes must be large enough to fill the
device) so we only need one C buffer */
devPtrC = (double *)(gpu_p->d_C);
/* ---------------------------------------------------------------------- */
/* copy Lx to the GPU */
/* ---------------------------------------------------------------------- */
/* copy host data to pinned buffer first for better H2D bandwidth */
#pragma omp parallel for num_threads(CHOLMOD_OMP_NUM_THREADS) if (ndcol > 32)
for ( icol=0; icol<ndcol; icol++ ) {
for ( irow=0; irow<ndrow2*L_ENTRY; irow++ ) {
gpu_p->h_Lx[iHostBuff][icol*ndrow2*L_ENTRY+irow] =
Lx[pdx1*L_ENTRY+icol*ndrow*L_ENTRY + irow];
}
}
cudaStat[0] = cudaMemcpyAsync ( devPtrLx,
gpu_p->h_Lx[iHostBuff],
ndrow2*ndcol*L_ENTRY*sizeof(devPtrLx[0]),
cudaMemcpyHostToDevice,
Common->gpuStream[iDevBuff] );
if ( cudaStat[0] ) {
CHOLMOD_GPU_PRINTF ((" ERROR cudaMemcpyAsync = %d \n", cudaStat[0]));
return (0);
}
/* make the current stream wait for kernels in previous streams */
cudaStreamWaitEvent ( Common->gpuStream[iDevBuff],
Common->updateCKernelsComplete, 0 ) ;
/* ---------------------------------------------------------------------- */
/* create the relative map for this descendant supernode */
/* ---------------------------------------------------------------------- */
createRelativeMapOnDevice ( (Int *)(gpu_p->d_Map),
(Int *)(gpu_p->d_Ls),
(Int *)(gpu_p->d_RelativeMap),
pdi1, ndrow2,
&(Common->gpuStream[iDevBuff]) );
/* ---------------------------------------------------------------------- */
/* do the CUDA SYRK */
/* ---------------------------------------------------------------------- */
cublasStatus = cublasSetStream (Common->cublasHandle,
Common->gpuStream[iDevBuff]) ;
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS stream") ;
}
alpha = 1.0 ;
beta = 0.0 ;
#ifdef REAL
cublasStatus = cublasDsyrk (Common->cublasHandle,
CUBLAS_FILL_MODE_LOWER,
CUBLAS_OP_N,
(int) ndrow1,
(int) ndcol, /* N, K: L1 is ndrow1-by-ndcol */
&alpha, /* ALPHA: 1 */
devPtrLx,
ndrow2, /* A, LDA: L1, ndrow2 */
&beta, /* BETA: 0 */
devPtrC,
ndrow2) ; /* C, LDC: C1 */
#else
cublasStatus = cublasZherk (Common->cublasHandle,
CUBLAS_FILL_MODE_LOWER,
CUBLAS_OP_N,
(int) ndrow1,
(int) ndcol, /* N, K: L1 is ndrow1-by-ndcol*/
&alpha, /* ALPHA: 1 */
(const cuDoubleComplex *) devPtrLx,
ndrow2, /* A, LDA: L1, ndrow2 */
&beta, /* BETA: 0 */
(cuDoubleComplex *) devPtrC,
ndrow2) ; /* C, LDC: C1 */
#endif
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS routine failure") ;
}
#ifndef NTIMER
Common->CHOLMOD_GPU_SYRK_TIME += SuiteSparse_time() - Common->syrkStart;
#endif
/* ---------------------------------------------------------------------- */
/* compute remaining (ndrow2-ndrow1)-by-ndrow1 block of C, C2 = L2*L1' */
/* ---------------------------------------------------------------------- */
#ifndef NTIMER
Common->CHOLMOD_GPU_GEMM_CALLS++ ;
tstart = SuiteSparse_time();
#endif
if (ndrow3 > 0)
{
#ifndef REAL
cuDoubleComplex calpha = {1.0,0.0} ;
cuDoubleComplex cbeta = {0.0,0.0} ;
#endif
/* ------------------------------------------------------------------ */
/* do the CUDA BLAS dgemm */
/* ------------------------------------------------------------------ */
#ifdef REAL
alpha = 1.0 ;
beta = 0.0 ;
cublasStatus = cublasDgemm (Common->cublasHandle,
CUBLAS_OP_N, CUBLAS_OP_T,
ndrow3, ndrow1, ndcol, /* M, N, K */
&alpha, /* ALPHA: 1 */
devPtrLx + L_ENTRY*(ndrow1), /* A, LDA: L2*/
ndrow2, /* ndrow */
devPtrLx, /* B, LDB: L1 */
ndrow2, /* ndrow */
&beta, /* BETA: 0 */
devPtrC + L_ENTRY*ndrow1, /* C, LDC: C2 */
ndrow2) ;
#else
cublasStatus = cublasZgemm (Common->cublasHandle,
CUBLAS_OP_N, CUBLAS_OP_C,
ndrow3, ndrow1, ndcol, /* M, N, K */
&calpha, /* ALPHA: 1 */
(const cuDoubleComplex*) devPtrLx + ndrow1,
ndrow2, /* ndrow */
(const cuDoubleComplex *) devPtrLx,
ndrow2, /* ndrow */
&cbeta, /* BETA: 0 */
(cuDoubleComplex *)devPtrC + ndrow1,
ndrow2) ;
#endif
if (cublasStatus != CUBLAS_STATUS_SUCCESS)
{
ERROR (CHOLMOD_GPU_PROBLEM, "GPU CUBLAS routine failure") ;
}
}
#ifndef NTIMER
Common->CHOLMOD_GPU_GEMM_TIME += SuiteSparse_time() - tstart;
#endif
/* ------------------------------------------------------------------ */
/* Assemble the update C on the device using the d_RelativeMap */
/* ------------------------------------------------------------------ */
#ifdef REAL
addUpdateOnDevice ( gpu_p->d_A[0], devPtrC,
gpu_p->d_RelativeMap, ndrow1, ndrow2, nsrow,
&(Common->gpuStream[iDevBuff]) );
#else
addComplexUpdateOnDevice ( gpu_p->d_A[0], devPtrC,
gpu_p->d_RelativeMap, ndrow1, ndrow2, nsrow,
&(Common->gpuStream[iDevBuff]) );
#endif
/* Record an event indicating that kernels for
this descendant are complete */
cudaEventRecord ( Common->updateCKernelsComplete,
Common->gpuStream[iDevBuff]);
cudaEventRecord ( Common->updateCBuffersFree[iHostBuff],
Common->gpuStream[iDevBuff]);
return (1) ;
}
void Stream::wait( EventSharedPtr const& event )
{
CUDA_VERIFY( cudaStreamWaitEvent( m_stream, event->getEvent(), 0 ) );
}
extern "C" magma_int_t
magma_ssytrd_sy2sb( char uplo, magma_int_t n, magma_int_t nb,
float *a, magma_int_t lda,
float *tau,
float *work, magma_int_t lwork,
float *dT,
magma_int_t threads, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SSYTRD_HE2HB reduces a real symmetric matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**T * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = 'U', the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
dT (output) REAL array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a + ((a_2)-1)*( lda) + (a_1)-1)
#define da_ref(a_1,a_2) (da + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define t_ref(a_1) (dT + ((a_1)-1)*(lddt))
char uplo_[2] = {uplo, 0};
int ldda = ((n+31)/32)*32;
int lddt = nb;
float c_neg_one = MAGMA_S_NEG_ONE;
float c_neg_half = MAGMA_S_NEG_HALF;
float c_one = MAGMA_S_ONE ;
float c_zero = MAGMA_S_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0;
int i;
int lwkopt;
int lquery;
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
if (*info == 0) {
/* Determine the block size. */
lwkopt = n * nb;
MAGMA_S_SET2REAL( work[0], lwkopt );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
float *da;
if (MAGMA_SUCCESS != magma_smalloc( &da, (n + 2*nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_int_t mklth = min(threads,12);
#if defined(USEMKL)
mkl_set_num_threads(mklth);
#endif
#if defined(USEACML)
omp_set_num_threads(mklth);
#endif
/* Use the first panel of da as work space */
float *dwork = da+n*ldda;
float *dW = dwork + nb*ldda;
#ifdef TRACING
char buf[80];
#endif
cudaStream_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
stream[2] = 0; // default stream
trace_init( 1, 1, 3, stream );
float *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(float));
magmablasSetKernelStream( stream[0] );
cudaEvent_t Pupdate_event;
cudaEventCreateWithFlags(&Pupdate_event,cudaEventDisableTiming);
//cudaEventCreate(&Pupdate_event);
if (upper) {
printf("SSYTRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
}else {
/* Copy the matrix to the GPU */
if (1 <= n-nb){
trace_gpu_start( 0, 0, "set", "set A" );
magma_ssetmatrix_async( (n-nb), (n-nb),
a_ref(nb+1, nb+1), lda,
da_ref(nb+1, nb+1), ldda, stream[0] );
trace_gpu_end( 0, 0 );
}
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb)
{
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ){
// spanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
spanel_to_q(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
trace_gpu_start( 0, 1, "get", "get panel" );
//magma_queue_sync( stream[0] );
cudaStreamWaitEvent(stream[1], Pupdate_event, 0);
magma_sgetmatrix_async( (pm+pn), pn,
da_ref( i, i), ldda,
a_ref ( i, i), lda, stream[1] );
trace_gpu_end( 0, 1 );
trace_gpu_start( 0, 2, "syr2k", "syr2k" );
magma_ssyr2k(MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old, c_neg_one,
da_ref(indi_old+pn_old, indj_old), ldda,
dW + pn_old , pm_old, d_one,
da_ref(indi_old+pn_old, indi_old+pn_old), ldda);
trace_gpu_end( 0, 2 );
trace_cpu_start( 0, "sync", "sync on 1" );
magma_queue_sync( stream[1] );
trace_cpu_end( 0 );
sq_to_panel(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
#ifdef TRACING
snprintf( buf, sizeof(buf), "panel %d", i );
#endif
trace_cpu_start( 0, "geqrf", buf );
lapackf77_sgeqrf(&pm, &pn, a_ref(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, a_ref(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
spanel_to_q(MagmaUpper, pk, a_ref(indi, indj), lda, work);
trace_cpu_end( 0 );
/* Send V from the CPU to the GPU */
trace_gpu_start( 0, 0, "set", "set V and T" );
magma_ssetmatrix_async( pm, pk,
a_ref(indi, indj), lda,
da_ref(indi, indj), ldda, stream[0] );
/* Send the triangular factor T to the GPU */
magma_ssetmatrix_async( pk, pk,
hT, nb,
t_ref(i), lddt, stream[0] );
trace_gpu_end( 0, 0 );
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
/* dwork = V T */
trace_cpu_start( 0, "sync", "sync on 0" );
// this sync is done here to be sure that the copy has been finished
// because below we made a restore sq_to_panel and this restore need
// to ensure that the copy has been finished. we did it here to allow
// overlapp of restore with next gemm and symm.
magma_queue_sync( stream[0] );
trace_cpu_end( 0 );
trace_gpu_start( 0, 2, "gemm", "work = V*T" );
magma_sgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, da_ref(indi, indj), ldda,
t_ref(i), lddt,
c_zero, dwork, pm);
trace_gpu_end( 0, 2 );
/* dW = X = A*V*T. dW = A*dwork */
trace_gpu_start( 0, 2, "symm", "X = A*work" );
magma_ssymm(MagmaLeft, uplo, pm, pk,
c_one, da_ref(indi, indi), ldda,
dwork, pm,
c_zero, dW, pm);
trace_gpu_end( 0, 2 );
/* restore the panel */
sq_to_panel(MagmaUpper, pk, a_ref(indi, indj), lda, work);
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
trace_gpu_start( 0, 2, "gemm", "work = T'*V'*X" );
magma_sgemm(MagmaTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork, pm,
dW, pm,
c_zero, dwork + pm*nb, nb);
trace_gpu_end( 0, 2 );
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
trace_gpu_start( 0, 2, "gemm", "W = X - 0.5*V*(T'*V'*X)" );
magma_sgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, da_ref(indi, indj), ldda,
dwork + pm*nb, nb,
c_one, dW, pm);
trace_gpu_end( 0, 2 );
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb){
/* There would be next iteration;
do lookahead - update the next panel */
trace_gpu_start( 0, 2, "gemm", "gemm 4 next panel left" );
magma_sgemm(MagmaNoTrans, MagmaTrans, pm, pn, pn, c_neg_one,
da_ref(indi, indj), ldda,
dW , pm, c_one,
da_ref(indi, indi), ldda);
trace_gpu_end( 0, 2 );
trace_gpu_start( 0, 2, "gemm", "gemm 5 next panel right" );
magma_sgemm(MagmaNoTrans, MagmaTrans, pm, pn, pn, c_neg_one,
dW , pm,
da_ref(indi, indj), ldda, c_one,
da_ref(indi, indi), ldda);
trace_gpu_end( 0, 2 );
cudaEventRecord(Pupdate_event, stream[0]);
}
else {
/* no look-ahead as this is last iteration */
trace_gpu_start( 0, 2, "syr2k", "syr2k last iteration" );
magma_ssyr2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
da_ref(indi, indj), ldda,
dW , pm, d_one,
da_ref(indi, indi), ldda);
trace_gpu_end( 0, 2 );
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for(i)
/* Send the last block to the CPU */
pk = min(pm,pn);
if (1 <= n-nb){
spanel_to_q(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
trace_gpu_start( 0, 2, "get", "get last block" );
magma_sgetmatrix( pk, pk,
da_ref(n-pk+1, n-pk+1), ldda,
a_ref(n-pk+1, n-pk+1), lda );
trace_gpu_end( 0, 2 );
sq_to_panel(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
}
}// end of LOWER
trace_finalize( "ssytrd_sy2sb.svg", "trace.css" );
cudaEventDestroy(Pupdate_event);
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( da );
MAGMA_S_SET2REAL( work[0], lwkopt );
magmablasSetKernelStream( 0 );
#if defined(USEMKL)
mkl_set_num_threads(1);
#endif
#if defined(USEACML)
omp_set_num_threads(1);
#endif
return *info;
} /* ssytrd_sy2sb_ */
/**
Purpose
-------
CHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = MagmaUpper, the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT COMPLEX array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_cheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_chetrd_he2hb( magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magmaFloatComplex *dT,
magma_int_t *info)
{
#define A(a_1,a_2) ( A + ((a_2)-1)*( lda) + (a_1)-1)
#define dA(a_1,a_2) (dA + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define dT(a_1) (dT + ((a_1)-1)*(lddt))
int ldda = ((n+31)/32)*32;
int lddt = nb;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0;
int i;
int lwkopt;
int lquery;
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magmaFloatComplex *dA;
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (n + 2*nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_int_t threads = magma_get_lapack_numthreads();
magma_int_t mklth = min(threads,16);
magma_set_lapack_numthreads(mklth);
/* Use the first panel of dA as work space */
magmaFloatComplex *dwork = dA + n*ldda;
magmaFloatComplex *dW = dwork + nb*ldda;
#ifdef TRACING
char buf[80];
#endif
magma_queue_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
stream[2] = 0; // default stream
trace_init( 1, 1, 3, stream );
magmaFloatComplex *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(magmaFloatComplex));
magmablasSetKernelStream( stream[0] );
magma_event_t Pupdate_event;
cudaEventCreateWithFlags(&Pupdate_event,cudaEventDisableTiming);
//cudaEventCreate(&Pupdate_event);
if (upper) {
printf("CHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
} else {
/* Copy the matrix to the GPU */
if (1 <= n-nb) {
trace_gpu_start( 0, 0, "set", "set A" );
magma_csetmatrix_async( (n-nb), (n-nb),
A(nb+1, nb+1), lda,
dA(nb+1, nb+1), ldda, stream[0] );
trace_gpu_end( 0, 0 );
}
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb) {
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ) {
// cpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
cpanel_to_q(MagmaUpper, pn-1, A(i, i+1), lda, work);
trace_gpu_start( 0, 1, "get", "get panel" );
//magma_queue_sync( stream[0] );
cudaStreamWaitEvent(stream[1], Pupdate_event, 0);
magma_cgetmatrix_async( (pm+pn), pn,
dA( i, i), ldda,
A ( i, i), lda, stream[1] );
trace_gpu_end( 0, 1 );
trace_gpu_start( 0, 2, "her2k", "her2k" );
magma_cher2k(MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old, c_neg_one,
dA(indi_old+pn_old, indj_old), ldda,
dW + pn_old, pm_old, d_one,
dA(indi_old+pn_old, indi_old+pn_old), ldda);
trace_gpu_end( 0, 2 );
trace_cpu_start( 0, "sync", "sync on 1" );
magma_queue_sync( stream[1] );
trace_cpu_end( 0 );
cq_to_panel(MagmaUpper, pn-1, A(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
#ifdef TRACING
snprintf( buf, sizeof(buf), "panel %d", i );
#endif
trace_cpu_start( 0, "geqrf", buf );
lapackf77_cgeqrf(&pm, &pn, A(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, A(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
cpanel_to_q(MagmaUpper, pk, A(indi, indj), lda, work);
trace_cpu_end( 0 );
/* Send V from the CPU to the GPU */
trace_gpu_start( 0, 0, "set", "set V and T" );
magma_csetmatrix_async( pm, pk,
A(indi, indj), lda,
dA(indi, indj), ldda, stream[0] );
/* Send the triangular factor T to the GPU */
magma_csetmatrix_async( pk, pk,
hT, nb,
dT(i), lddt, stream[0] );
trace_gpu_end( 0, 0 );
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
/* dwork = V T */
trace_cpu_start( 0, "sync", "sync on 0" );
// this sync is done here to be sure that the copy has been finished
// because below we made a restore cq_to_panel and this restore need
// to ensure that the copy has been finished. we did it here to allow
// overlapp of restore with next gemm and symm.
magma_queue_sync( stream[0] );
trace_cpu_end( 0 );
trace_gpu_start( 0, 2, "gemm", "work = V*T" );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dA(indi, indj), ldda,
dT(i), lddt,
c_zero, dwork, pm);
trace_gpu_end( 0, 2 );
/* dW = X = A*V*T. dW = A*dwork */
trace_gpu_start( 0, 2, "hemm", "X = A*work" );
magma_chemm(MagmaLeft, uplo, pm, pk,
c_one, dA(indi, indi), ldda,
dwork, pm,
c_zero, dW, pm);
trace_gpu_end( 0, 2 );
/* restore the panel */
cq_to_panel(MagmaUpper, pk, A(indi, indj), lda, work);
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
trace_gpu_start( 0, 2, "gemm", "work = T'*V'*X" );
magma_cgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork, pm,
dW, pm,
c_zero, dwork + pm*nb, nb);
trace_gpu_end( 0, 2 );
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
trace_gpu_start( 0, 2, "gemm", "W = X - 0.5*V*(T'*V'*X)" );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dA(indi, indj), ldda,
dwork + pm*nb, nb,
c_one, dW, pm);
trace_gpu_end( 0, 2 );
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb) {
/* There would be next iteration;
do lookahead - update the next panel */
trace_gpu_start( 0, 2, "gemm", "gemm 4 next panel left" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dA(indi, indj), ldda,
dW, pm, c_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
trace_gpu_start( 0, 2, "gemm", "gemm 5 next panel right" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dW, pm,
dA(indi, indj), ldda, c_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
cudaEventRecord(Pupdate_event, stream[0]);
}
else {
/* no look-ahead as this is last iteration */
trace_gpu_start( 0, 2, "her2k", "her2k last iteration" );
magma_cher2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dA(indi, indj), ldda,
dW, pm, d_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for (i)
/* Send the last block to the CPU */
pk = min(pm,pn);
if (1 <= n-nb) {
cpanel_to_q(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
trace_gpu_start( 0, 2, "get", "get last block" );
magma_cgetmatrix( pk, pk,
dA(n-pk+1, n-pk+1), ldda,
A(n-pk+1, n-pk+1), lda );
trace_gpu_end( 0, 2 );
cq_to_panel(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
}
}// end of LOWER
trace_finalize( "chetrd_he2hb.svg", "trace.css" );
cudaEventDestroy(Pupdate_event);
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dA );
work[0] = MAGMA_C_MAKE( lwkopt, 0 );
magmablasSetKernelStream( 0 );
magma_set_lapack_numthreads(threads);
return *info;
} /* magma_chetrd_he2hb */
