Merge branch 'charm' of charmgit:charm into charm
authorGengbin Zheng <tg455581@login3.ranger.tacc.utexas.edu>
Sat, 20 Feb 2010 19:50:08 +0000 (13:50 -0600)
committerGengbin Zheng <tg455581@login3.ranger.tacc.utexas.edu>
Sat, 20 Feb 2010 19:50:08 +0000 (13:50 -0600)
src/arch/cuda/hybridAPI/cuda-hybrid-api.cu
src/arch/cuda/hybridAPI/cuda-hybrid-api.h
src/arch/cuda/hybridAPI/wr.h
src/ck-core/ckmemcheckpoint.C
src/conv-core/convcore.c
src/libs/ck-libs/pose/rep.h

index bc188cc132c3b898b747e66d07d14e8b907c2184..03994d4e3ae58f0b1610a2de4f50259bddbfe913 100644 (file)
  *  and executes the callback 
  */ 
 extern void CUDACallbackManager(void * fn); 
+extern int CmiMyPe();
 
 /* initial size of the user-addressed portion of host/device buffer
  * arrays; the system-addressed portion of host/device buffer arrays
  * (used when there is no need to share buffers between work requests)
  * will be equivalant in size.  
  */ 
-#define NUM_BUFFERS 128
+#define NUM_BUFFERS 256
 #define MAX_PINNED_REQ 64  
 #define MAX_DELAYED_FREE_REQS 64  
 
@@ -52,6 +53,13 @@ pinnedMemReq pinnedMemQueue[MAX_PINNED_REQ];
 unsigned int currentDfr = 0;
 DelayedFreeReq delayedFreeReqs[MAX_DELAYED_FREE_REQS];
 
+#ifdef GPU_MEMPOOL
+#define GPU_MEMPOOL_NUM_SLOTS 15
+
+CkVec<BufferPool> memPoolFreeBufs;
+CkVec<int> memPoolBoundaries;
+//int memPoolBoundaries[GPU_MEMPOOL_NUM_SLOTS];
+#endif
 
 /* The runtime system keeps track of all allocated buffers on the GPU.
  * The following arrays contain pointers to host (CPU) data and the
@@ -270,8 +278,9 @@ void allocateBuffers(workRequest *wr) {
       // allocate if the buffer for the corresponding index is NULL 
       if (devBuffers[index] == NULL && size > 0) {
 #ifdef GPU_PRINT_BUFFER_ALLOCATE
-        double mil = 1e6;
-        printf("*** ALLOCATE buffer 0x%x (%d) size %f mb\n", devBuffers[index], index, 1.0*size/mil);
+        double mil = 1e3;
+        printf("*** ALLOCATE buffer 0x%x (%d) size %f kb\n", devBuffers[index], index, 1.0*size/mil);
+
 #endif
 
         CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &devBuffers[index], size));
@@ -433,6 +442,31 @@ void initHybridAPI(int myPe) {
   traceRegisterUserEvent("GPU Memory Cleanup", GPU_MEM_CLEANUP);
 #endif
 
+#ifdef GPU_MEMPOOL
+  int nslots = GPU_MEMPOOL_NUM_SLOTS;
+  int *sizes;
+  sizes = (int *)malloc(sizeof(int)*nslots); 
+
+  memPoolBoundaries.reserve(GPU_MEMPOOL_NUM_SLOTS);
+  memPoolBoundaries.length() = GPU_MEMPOOL_NUM_SLOTS;
+
+  int bufSize = GPU_MEMPOOL_MIN_BUFFER_SIZE;
+  for(int i = 0; i < GPU_MEMPOOL_NUM_SLOTS; i++){
+    memPoolBoundaries[i] = bufSize;
+    bufSize = bufSize << 1;
+  }
+
+  sizes[0] = sizes[1] = sizes[2] = sizes[3] = 32; 
+  sizes[4] = sizes[5] = sizes[6] = sizes[7] = 16; 
+  sizes[8] = sizes[9] = sizes[10] = sizes[11] = 4; 
+  sizes[12] = sizes[13] = sizes[14] = 2; 
+
+  printf("creating buffer pool...");
+  createPool(sizes, nslots, memPoolFreeBufs);
+  printf("...done\n");
+
+#endif
+
   currentDfr = 0;
   for(int i = 0; i < MAX_DELAYED_FREE_REQS; i++){
     delayedFreeReqs[i].freed = true;
@@ -623,7 +657,7 @@ void gpuProgressFn() {
       */
   }
   if (head->state == TRANSFERRING_OUT) {
-    if (cudaStreamQuery(data_out_stream) == cudaSuccess{
+    if (cudaStreamQuery(data_out_stream) == cudaSuccess && cudaStreamQuery(kernel_stream) == cudaSuccess){
       freeMemory(head); 
 #ifdef GPU_PROFILE
       gpuEvents[dataCleanupIndex].endTime = cutGetTimerValue(timerHandle);
@@ -640,6 +674,9 @@ void gpuProgressFn() {
   }
 }
 
+#ifdef GPU_MEMPOOL
+void releasePool(CkVec<BufferPool> &pools);
+#endif
 /* exitHybridAPI
  *  cleans up and deletes memory allocated for the queue and the CUDA streams
  */
@@ -671,4 +708,115 @@ void exitHybridAPI() {
   CUT_SAFE_CALL(cutDeleteTimer(timerHandle));  
 
 #endif
+
+#ifdef GPU_MEMPOOL
+  releasePool(memPoolFreeBufs);
+#endif
+
+}
+
+#ifdef GPU_MEMPOOL
+void releasePool(CkVec<BufferPool> &pools){
+  for(int i = 0; i < pools.length(); i++){
+    CUDA_SAFE_CALL_NO_SYNC(cudaFreeHost((void *)pools[i].head));
+  }
+  pools.free();
+}
+
+// Create a pool with nslots slots.
+// There are nbuffers[i] buffers for each buffer size corresponding to entry i
+// FIXME - list the alignment/fragmentation issues with either of two allocation schemes:
+// if a single, large buffer is allocated for each subpool
+// if multiple smaller buffers are allocated for each subpool
+void createPool(int *nbuffers, int nslots, CkVec<BufferPool> &pools){
+  //pools  = (BufferPool *)malloc(nslots*sizeof(BufferPool));
+  pools.reserve(nslots);
+  pools.length() = nslots;
+
+  for(int i = 0; i < nslots; i++){
+    int bufSize = memPoolBoundaries[i];
+    int numBuffers = nbuffers[i];
+    pools[i].size = bufSize;
+    
+    CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)(&pools[i].head), 
+                                          (sizeof(Header)+bufSize)*numBuffers));
+    if(pools[i].head == NULL){
+      abort();
+    }
+
+    Header *hd = pools[i].head;
+    Header *previous = NULL;
+    char *memory;
+
+    for(int j = 0; j < numBuffers; j++){
+      hd->slot = i;
+      hd->next = previous;
+      previous = hd;
+      hd++; // move ptr past header
+      memory = (char *)hd;
+      memory += bufSize;
+      hd = (Header *)memory;
+    }
+
+    pools[i].head = previous;
+  }
+}
+
+int findPool(int size){
+  int boundaryArrayLen = memPoolBoundaries.length();
+  if(size <= memPoolBoundaries[0]){
+    return (0);
+  }
+  else if(size > memPoolBoundaries[boundaryArrayLen-1]){
+    // create new slot
+    memPoolBoundaries.push_back(size);
+
+    BufferPool newpool;
+    CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)&newpool.head, size+sizeof(Header)));
+    newpool.size = size;
+    memPoolFreeBufs.push_back(newpool);
+
+    Header *hd = newpool.head;
+    hd->next = NULL;
+    hd->slot = boundaryArrayLen;
+
+    return boundaryArrayLen;
+  }
+  for(int i = 0; i < GPU_MEMPOOL_NUM_SLOTS-1; i++){
+    if(memPoolBoundaries[i] < size && size <= memPoolBoundaries[i+1]){
+      return (i+1);
+    }
+  }
+  return -1;
+}
+
+void *getBufferFromPool(int pool, int size){
+  Header *ret;
+  if(pool < 0 || pool >= memPoolFreeBufs.length() || memPoolFreeBufs[pool].head == NULL){
+    printf("(%d) pool %d size: %d\n", CmiMyPe(), pool, size);
+    abort();
+  }
+  else{
+    ret = memPoolFreeBufs[pool].head;
+    memPoolFreeBufs[pool].head = ret->next;
+    return (void *)(ret+1);
+  }
+  return NULL;
+}
+
+void returnBufferToPool(int pool, Header *hd){
+  hd->next = memPoolFreeBufs[pool].head;
+  memPoolFreeBufs[pool].head = hd;
+}
+
+void *hapi_poolMalloc(int size){
+  return getBufferFromPool(findPool(size), size);
+}
+
+void hapi_poolFree(void *ptr){
+  Header *hd = ((Header *)ptr)-1;
+  returnBufferToPool(hd->slot, hd);
 }
+
+
+#endif
index 16482596140a587906ae0a6822360fe31daaea93..77cb08f8803dbe445ec47f87fd9edb4082547a35 100644 (file)
 #ifndef __CUDA_HYBRID_API_H__
 #define __CUDA_HYBRID_API_H__
 
+#ifdef GPU_MEMPOOL
+#include "cklists.h"
+#endif
+
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -36,6 +40,28 @@ void gpuProgressFn();
 */
 void exitHybridAPI(); 
 
+
+#ifdef GPU_MEMPOOL
+// data and metadata reside in same chunk of memory
+typedef struct _header{
+  //void *buf;
+  struct _header *next;
+  int slot;
+}Header;
+
+typedef struct _bufferPool{
+  Header *head;
+  //bool expanded;
+  int size;
+}BufferPool;
+
+// pre-allocated buffers will be at least this big
+#define GPU_MEMPOOL_MIN_BUFFER_SIZE 1024
+
+void createPool(int *nbuffers, int nslots, CkVec<BufferPool> &pools);
+
+#endif
+
 #ifdef __cplusplus
 }
 #endif
index a0316dd592076468d2d4d780df5f4d37e04ef909..9b688814bfa31ca625932b609be2fe03fbf967a9 100644 (file)
@@ -31,6 +31,7 @@ typedef struct _delayedFreeReq {
 
 void delayedFree(void *ptr);
 
+
 /* pinnedMallocHost
  *
  * schedules a pinned memory allocation so that it does not impede
@@ -163,6 +164,10 @@ typedef struct {
  */
 void enqueue(workRequestQueue *q, workRequest *wr); 
 
+#ifdef GPU_MEMPOOL
+void hapi_poolFree(void *);
+void *hapi_poolMalloc(int size);
+#endif
 /* external declarations needed by the user */
 
 extern workRequestQueue *wrQueue; 
index f8be9cbbb5a75ecb1c4ca007508414ea33a82976..29af2288e59132585e864fd09bcc6c7c07935ed5 100644 (file)
@@ -48,7 +48,7 @@ TODO:
 #include <signal.h>
 
 // pick buddy processor from a different physical node
-#define NODE_CHECKPOINT                        1
+#define NODE_CHECKPOINT                        0
 
 // assume NO extra processors--1
 // assume extra processors--0
@@ -106,8 +106,8 @@ inline int CkMemCheckPT::BuddyPE(int pe)
   }
 #else
   budpe = pe;
-  while (budpe == pe || isFailed(budPe)) 
-          budPe = (budPe+1)%CkNumPes();
+  while (budpe == pe || isFailed(budpe)) 
+          budpe = (budpe+1)%CkNumPes();
 #endif
   return budpe;
 }
@@ -117,6 +117,9 @@ inline int CkMemCheckPT::BuddyPE(int pe)
 #if CMK_MEM_CHECKPOINT
 void ArrayElement::init_checkpt() {
        if (_memChkptOn == 0) return;
+       if (CkInRestarting()) {
+         CkPrintf("[%d] Warning: init_checkpt called during restart, possible bug in migration constructor!\n");
+       }
        // only master init checkpoint
         if (thisArray->getLocMgr()->firstManager->mgr!=thisArray) return;
 
@@ -461,7 +464,7 @@ void CkMemCheckPT::recvProcData(CkProcCheckPTMessage *msg)
 {
   if (CpvAccess(procChkptBuf)) delete CpvAccess(procChkptBuf);
   CpvAccess(procChkptBuf) = msg;
-//CmiPrintf("[%d] CkMemCheckPT::recvProcData report to %d\n", CkMyPe(), msg->reportPe);
+  DEBUGF("[%d] CkMemCheckPT::recvProcData report to %d\n", CkMyPe(), msg->reportPe);
   thisProxy[msg->reportPe].cpFinish();
 }
 
@@ -532,7 +535,7 @@ void CkMemCheckPT::report()
     objsize += entry->getSize();
   }
   CmiAssert(CpvAccess(procChkptBuf));
-  CkPrintf("[%d] Checkpointed Object size: %d len: %d Processor data: %d\n", CkMyPe(), objsize, len, CpvAccess(procChkptBuf)->len);
+  CkPrintf("[%d] Checkpoint object size: %d len: %d Processor data: %d \n", CkMyPe(), objsize, len, CpvAccess(procChkptBuf)->len);
 }
 
 /*****************************************************************************
@@ -755,12 +758,14 @@ void CkMemCheckPT::recoverArrayElements()
     inmem_restore(msg);
     count ++;
   }
-//CkPrintf("[%d] recoverArrayElements restore %d objects\n", CkMyPe(), count);
+  DEBUGF("[%d] recoverArrayElements restore %d objects\n", CkMyPe(), count);
 
   if (CkMyPe() == 0)
     CkStartQD(CkCallback(CkIndex_CkMemCheckPT::finishUp(), thisProxy));
 }
 
+static double restartT;
+
 // on every processor
 // turn load balancer back on
 void CkMemCheckPT::finishUp()
@@ -772,9 +777,12 @@ void CkMemCheckPT::finishUp()
 
   if (CkMyPe() == 0)
   {
-       CkPrintf("[%d] CkMemCheckPT ----- %s in %f seconds\n",CkMyPe(), stage, CmiWallTimer()-startTime);
+       CkPrintf("[%d] CkMemCheckPT ----- %s in %f seconds, callback triggered\n",CkMyPe(), stage, CmiWallTimer()-startTime);
        CkStartQD(cpCallback);
   } 
+  if (CkMyPe() == thisFailedPe)
+       CkPrintf("[%d] Restart finished in %f seconds.\n", CkMyPe(), CkWallTimer()-restartT);
+
 #if CK_NO_PROC_POOL
 #if NODE_CHECKPOINT
   int numnodes = CmiNumPhysicalNodes();
@@ -949,6 +957,7 @@ static void askProcDataHandler(char *msg)
 void CkMemRestart(const char *dummy, CkArgMsg *args)
 {
 #if CMK_MEM_CHECKPOINT
+   restartT = CkWallTimer();
    _diePE = CkMyPe();
    CmiPrintf("[%d] I am restarting  cur_restart_phase:%d \n",CmiMyPe(), cur_restart_phase);
    CkMemCheckPT::startTime = CmiWallTimer();
index d9d112e803a7fad57b3f312e6158def462b2f261..22846ca83d0af25fdd8b2bee784b5843af5e8414 100644 (file)
@@ -803,10 +803,12 @@ void CmiTimerInit()
   struct rusage ru;
   CpvInitialize(double, inittime_virtual);
 
+#if ! CMK_MEM_CHECKPOINT
   /* try to synchronize calling barrier */
   CmiBarrier();
   CmiBarrier();
   CmiBarrier();
+#endif
 
   gettimeofday(&tv,0);
   inittime_wallclock = (tv.tv_sec * 1.0) + (tv.tv_usec*0.000001);
@@ -815,8 +817,10 @@ void CmiTimerInit()
     (ru.ru_utime.tv_sec * 1.0)+(ru.ru_utime.tv_usec * 0.000001) +
     (ru.ru_stime.tv_sec * 1.0)+(ru.ru_stime.tv_usec * 0.000001);
 
+#if ! CMK_MEM_CHECKPOINT
   CmiBarrier();
 /*  CmiBarrierZero(); */
+#endif
 }
 
 double CmiCpuTimer()
index 5e908ed35a77d047e715113edfb19a7cf063d6c7..fb597c0b5b3f4337b3da42fe10aefd527eba82c1 100644 (file)
@@ -127,8 +127,13 @@ class rep
     prand_seed = rnum+INT_MAX; 
     return prand_seed;
   }
+#if !defined(_WIN32) || defined(__CYGWIN__) 
   inline long int POSE_Linear_rand() { return nrand48(prand48_seed); }
   inline double POSE_Uniform_rand() { return erand48(prand48_seed); }
+#else
+  inline long int POSE_Linear_rand() { return CrnRand(); }
+  inline double POSE_Uniform_rand() { return 1.0*CrnRand()/MAXINT; }
+#endif
 
 };