Merge branch 'charm' of charmgit:charm into charm
authorFilippo Gioachin <gioachin@uiuc.edu>
Wed, 7 Jul 2010 09:10:36 +0000 (04:10 -0500)
committerFilippo Gioachin <gioachin@uiuc.edu>
Wed, 7 Jul 2010 09:10:36 +0000 (04:10 -0500)
.gitignore
examples/ParFUM/simple2D/Makefile
src/arch/cuda/hybridAPI/Makefile
src/arch/cuda/hybridAPI/cuda-hybrid-api.cu
src/arch/cuda/hybridAPI/cuda-hybrid-api.h
src/arch/net-linux-amd64-cuda/Makefile.machine [new file with mode: 0644]
src/arch/net-linux-amd64-cuda/special.sh
src/arch/net/charmrun/Makefile
src/libs/ck-libs/cache/Makefile
src/libs/ck-libs/liveViz/Makefile

index d4763e22636362a49a54859e8c1a5f8aaef039bb..0c74885e2b77422261eac21aa1b2dacb281f8bbd 100644 (file)
@@ -17,13 +17,13 @@ lib_so
 tmp
 
 # Ignore directories created from compilation
-net-*
-mpi-*
-lapi*
-vmi-*
-multicore-*
-uth-*
-bluegenep*
+/net-*
+/mpi-*
+/lapi*
+/vmi-*
+/multicore-*
+/uth-*
+/bluegenep*
 
 # Ignore cruft that would result from compiling charm example/test code
 # Note this is negated inside src/arch to permit addition of new charmrun files in there
index e6de92ba48f13965a0f742aa0707c6dc110ff3a0..87e77be4078a1702ac93fec0644bc37e88e2ac47 100644 (file)
@@ -1,4 +1,4 @@
-CHARMC=../../../bin/charmc $(OPTS) 
+CHARMC=../../../bin/charmc $(OPTS)
 
 all: pgm
 
index 4a4ad0c37dd50a447039b8c3f36668b1bfa5896d..639071004b1c11b00ea592586d92ed352cd31b6e 100644 (file)
@@ -1,8 +1,10 @@
+CUDA_DIR=/usr/local/cuda/cuda
+NVIDIA_CUDA_SDK=/usr/local/NVIDIA_CUDA_SDK
 CHARMDIR=../../
 CHARMC=$(CHARMDIR)/bin/charmc
 
 NVCC = $(CUDA_DIR)/bin/nvcc
-NVCC_FLAGS = -c -use_fast_math #-DGPU_MEMPOOL -DGPU_MEMPOOL_DEBUG -DGPU_WRQ_VERBOSE #-device-debug -deviceemu
+NVCC_FLAGS = -c -use_fast_math -DGPU_MEMPOOL #-DGPU_PROFILE -DGPU_TRACE #-DGPU_MEMPOOL_DEBUG -DGPU_WRQ_VERBOSE #-device-debug -deviceemu
 NVCC_INC = -I$(CUDA_DIR)/include -I$(NVIDIA_CUDA_SDK)/common/inc -I.. 
 
 RM = rm -f
index eec793c726bf0ce4aea1b264ddc982592dec6970..14c73dc0d22fad75bae52ea52aeeb0867c4d2712 100644 (file)
@@ -58,11 +58,19 @@ unsigned int currentDfr = 0;
 void *delayedFreeReqs[MAX_DELAYED_FREE_REQS];
 
 #ifdef GPU_MEMPOOL
-#define GPU_MEMPOOL_NUM_SLOTS 15
+#define GPU_MEMPOOL_NUM_SLOTS 19
+// pre-allocated buffers will be at least this big
+#define GPU_MEMPOOL_MIN_BUFFER_SIZE 256
 
 CkVec<BufferPool> memPoolFreeBufs;
 CkVec<int> memPoolBoundaries;
 //int memPoolBoundaries[GPU_MEMPOOL_NUM_SLOTS];
+
+#ifdef GPU_DUMMY_MEMPOOL
+CkVec<int> memPoolMax;
+CkVec<int> memPoolSize;
+#endif
+
 #endif
 
 /* The runtime system keeps track of all allocated buffers on the GPU.
@@ -448,9 +456,16 @@ void initHybridAPI(int myPe) {
 #endif
 
 #ifdef GPU_MEMPOOL
+
   int nslots = GPU_MEMPOOL_NUM_SLOTS;
-  int *sizes;
-  sizes = (int *)malloc(sizeof(int)*nslots); 
+  int sizes[GPU_MEMPOOL_NUM_SLOTS];
+
+#ifdef GPU_DUMMY_MEMPOOL
+  memPoolMax.reserve(nslots);
+  memPoolMax.length() = nslots;
+  memPoolSize.reserve(nslots);
+  memPoolSize.length() = nslots;
+#endif
 
   memPoolBoundaries.reserve(GPU_MEMPOOL_NUM_SLOTS);
   memPoolBoundaries.length() = GPU_MEMPOOL_NUM_SLOTS;
@@ -459,40 +474,37 @@ void initHybridAPI(int myPe) {
   for(int i = 0; i < GPU_MEMPOOL_NUM_SLOTS; i++){
     memPoolBoundaries[i] = bufSize;
     bufSize = bufSize << 1;
+#ifdef GPU_DUMMY_MEMPOOL
+    memPoolSize[i] = 0;
+    memPoolMax[i] = -1;
+#endif
   }
 
-  //1K
-  sizes[0] = 512; 
-  //2K
-  sizes[1] = 512;
-  //4K
-  sizes[2] = 64;
-  //8K
-  sizes[3] = 64;
-  //16K
-  sizes[4] = 32;
-  //32K
-  sizes[5] = 32;
-  //64K
-  sizes[6] = 32;
-  //128K
-  sizes[7] = 32;
-  //256K
-  sizes[8] = 32;
-  //512K
-  sizes[9] = 32;
-  //1M
-  sizes[10] = 170;
-  //2M
-  sizes[11] = 16;
-  //4M
-  sizes[12] = 4;
-  //8M
-  sizes[13] = 2;
-  //16M
-  sizes[14] = 2; 
-
-  createPool(sizes, nslots, memPoolFreeBufs);
+
+#ifndef GPU_DUMMY_MEMPOOL
+/*256*/ sizes[0] = 20;
+/*512*/ sizes[1] = 10;
+/*1024*/ sizes[2] = 10;
+/*2048*/ sizes[3] = 20;
+/*4096*/ sizes[4] = 10;
+/*8192*/ sizes[5] = 30;
+/*16384*/ sizes[6] = 25;
+/*32768*/ sizes[7] = 10;
+/*65536*/ sizes[8] = 5;
+/*131072*/ sizes[9] = 5;
+/*262144*/ sizes[10] = 5;
+/*524288*/ sizes[11] = 5;
+/*1048576*/ sizes[12] = 5;
+/*2097152*/ sizes[13] = 10;
+/*4194304*/ sizes[14] = 10;
+/*8388608*/ sizes[15] = 10;
+/*16777216*/ sizes[16] = 8;
+/*33554432*/ sizes[17] = 6;
+/*67108864*/ sizes[18] = 7;
+
+createPool(sizes, nslots, memPoolFreeBufs);
+#endif
+
   printf("[%d] done creating buffer pool\n", CmiMyPe());
 
 #endif
@@ -814,6 +826,24 @@ void releasePool(CkVec<BufferPool> &pools);
  *  cleans up and deletes memory allocated for the queue and the CUDA streams
  */
 void exitHybridAPI() {
+  printf("EXIT HYBRID API\n");
+
+#ifdef GPU_MEMPOOL
+
+#ifndef GPU_DUMMY_MEMPOOL
+  releasePool(memPoolFreeBufs);
+#else
+  for(int i = 0; i < memPoolBoundaries.length(); i++){
+    printf("(%d) slot %d size: %d max: %d\n", CmiMyPe(), i, memPoolBoundaries[i], memPoolMax[i]);
+  }
+
+  if(memPoolBoundaries.length() != memPoolMax.length()){
+    abort();
+  }
+#endif
+  
+#endif
+
   deleteWRqueue(wrQueue); 
   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(kernel_stream)); 
   CUDA_SAFE_CALL_NO_SYNC(cudaStreamDestroy(data_in_stream)); 
@@ -842,16 +872,18 @@ void exitHybridAPI() {
 
 #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));
+    Header *hdr;
+    Header *next;
+    for(hdr = pools[i].head; hdr != NULL;){
+      next = hdr->next; 
+      CUDA_SAFE_CALL_NO_SYNC(cudaFreeHost((void *)hdr));
+      hdr = next;
+    }
   }
   pools.free();
 }
@@ -870,25 +902,26 @@ void createPool(int *nbuffers, int nslots, CkVec<BufferPool> &pools){
     int bufSize = memPoolBoundaries[i];
     int numBuffers = nbuffers[i];
     pools[i].size = bufSize;
+    pools[i].head = NULL;
     
+    /*
     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++){
+      CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)&hd,
+                                            (sizeof(Header)+bufSize)));
+      if(hd == NULL){
+        printf("(%d) failed to allocate %dth block of size %d, slot %d\n", CmiMyPe(), j, bufSize, i);
+        abort();
+      }
       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;
@@ -906,9 +939,17 @@ int findPool(int size){
   else if(size > memPoolBoundaries[boundaryArrayLen-1]){
     // create new slot
     memPoolBoundaries.push_back(size);
+#ifdef GPU_DUMMY_MEMPOOL
+    memPoolMax.push_back(-1);
+    memPoolSize.push_back(0);
+#endif
 
     BufferPool newpool;
     CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)&newpool.head, size+sizeof(Header)));
+    if(newpool.head == NULL){
+      printf("(%d) findPool: failed to allocate newpool %d head, size %d\n", CmiMyPe(), boundaryArrayLen, size);
+      abort();
+    }
     newpool.size = size;
 #ifdef GPU_MEMPOOL_DEBUG
     newpool.num = 1;
@@ -921,7 +962,7 @@ int findPool(int size){
 
     return boundaryArrayLen;
   }
-  for(int i = 0; i < GPU_MEMPOOL_NUM_SLOTS-1; i++){
+  for(int i = 0; i < memPoolBoundaries.length()-1; i++){
     if(memPoolBoundaries[i] < size && size <= memPoolBoundaries[i+1]){
       return (i+1);
     }
@@ -931,12 +972,23 @@ int findPool(int size){
 
 void *getBufferFromPool(int pool, int size){
   Header *ret;
-  if(pool < 0 || pool >= memPoolFreeBufs.length() || memPoolFreeBufs[pool].head == NULL){
+  if(pool < 0 || pool >= memPoolFreeBufs.length()){
+    printf("(%d) getBufferFromPool, pool: %d, size: %d invalid pool\n", CmiMyPe(), pool, size);
 #ifdef GPU_MEMPOOL_DEBUG
-    printf("(%d) pool %d size: %d, num: %d\n", CmiMyPe(), pool, size, memPoolFreeBufs[pool].num);
+    printf("(%d) num: %d\n", CmiMyPe(), memPoolFreeBufs[pool].num);
 #endif
     abort();
   }
+  else if (memPoolFreeBufs[pool].head == NULL){
+    Header *hd;
+    CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)&hd, sizeof(Header)+memPoolFreeBufs[pool].size));
+    printf("(%d) getBufferFromPool, pool: %d, size: %d expand by 1\n", CmiMyPe(), pool, size);
+    if(hd == NULL){
+      abort();
+    }
+    hd->slot = pool;
+    return (void *)(hd+1);
+  }
   else{
     ret = memPoolFreeBufs[pool].head;
     memPoolFreeBufs[pool].head = ret->next;
@@ -959,7 +1011,25 @@ void returnBufferToPool(int pool, Header *hd){
 
 void *hapi_poolMalloc(int size){
   int pool = findPool(size);
-  void *buf = getBufferFromPool(pool, size);
+  void *buf;
+#ifdef GPU_DUMMY_MEMPOOL
+  CUDA_SAFE_CALL_NO_SYNC(cudaMallocHost((void **)&buf, size+sizeof(Header)));
+  if(pool < 0 || pool >= memPoolSize.length()){
+    printf("(%d) need to create up to pool %d; malloc size: %d\n", CmiMyPe(), pool, size);
+    abort();
+  }
+  memPoolSize[pool]++;
+  if(memPoolSize[pool] > memPoolMax[pool]){
+    memPoolMax[pool] = memPoolSize[pool];
+  }
+  Header *hdr = (Header *)buf;
+  hdr->slot = pool;
+  hdr = hdr+1;
+  buf = (void *)hdr;
+#else
+  buf = getBufferFromPool(pool, size);
+#endif
+  
 #ifdef GPU_MEMPOOL_DEBUG
   printf("(%d) hapi_malloc size %d pool %d left %d\n", CmiMyPe(), size, pool, memPoolFreeBufs[pool].num);
 #endif
@@ -969,7 +1039,18 @@ void *hapi_poolMalloc(int size){
 void hapi_poolFree(void *ptr){
   Header *hd = ((Header *)ptr)-1;
   int pool = hd->slot;
+
+#ifdef GPU_DUMMY_MEMPOOL
+  if(pool < 0 || pool >= memPoolSize.length()){
+    printf("(%d) free pool %d\n", CmiMyPe(), pool);
+    abort();
+  }
+  memPoolSize[pool]--;
+  delayedFree((void *)hd); 
+#else
   returnBufferToPool(pool, hd);
+#endif
+
 #ifdef GPU_MEMPOOL_DEBUG
   int size = hd->size;
   printf("(%d) hapi_free size %d pool %d left %d\n", CmiMyPe(), size, pool, memPoolFreeBufs[pool].num);
index 4e46988c141072e8724ad17674ce695c14973979..e51429a9214028ccf71c9780aec16c5acb2a62b5 100644 (file)
@@ -58,8 +58,6 @@ typedef struct _bufferPool{
 #endif
 }BufferPool;
 
-// pre-allocated buffers will be at least this big
-#define GPU_MEMPOOL_MIN_BUFFER_SIZE 1024
 
 
 #endif
diff --git a/src/arch/net-linux-amd64-cuda/Makefile.machine b/src/arch/net-linux-amd64-cuda/Makefile.machine
new file mode 100644 (file)
index 0000000..53db1e4
--- /dev/null
@@ -0,0 +1,6 @@
+hybridAPI:
+       cd hybridAPI && make install
+
+charm++: hybridAPI
+
+.PHONY: hybridAPI
index 4b6660d4ec2ea4cfb28553ec6ef0b34c412f14de..8290d1ed181ade84322d1e42cc8057ca2b6181c7 100755 (executable)
@@ -22,5 +22,3 @@ export CHARMINC=../include
 #PPU_EMBEDSPU = $CMK_PPU_EMBEDSPU
 #SPERT_LIBS = $CMK_SPERT_LIBS
 #EOF
-
-cd hybridAPI && make install
index 9e99bb14d47f113db87d03a8bb4f18b0e36dac01..6f5e6c562c35f35335f28d7ac8684c31afdeca8f 100644 (file)
@@ -1,5 +1,5 @@
 BIN=../../bin
-CHARMC=$(BIN)/charmc $(OPTS)
+CHARMC=$(BIN)/charmc $(OPTS) -I..
 
 SHELL=/bin/sh
 
index 90b056bedb0bac4dee2147735138f70e781136e4..7a14076295734fdcffce1aa467a4f8962c469143 100644 (file)
@@ -22,10 +22,10 @@ INTERFACE: $(CIFILES)
        touch INTERFACE
 
 CkCache.o: CkCache.C $(HEADERS)
-       $(CHARMC) -c -o CkCache.o CkCache.C
+       $(CHARMC) -I../../.. -c -o CkCache.o CkCache.C
 
 clean:
        rm -f conv-host *.o *.decl.h *.def.h core  $(LIB) INTERFACE
 
 realclean: clean
-       rm -f $(LIBDEST) $(HEADERS)
\ No newline at end of file
+       rm -f $(LIBDEST) $(HEADERS)
index ad146ac9076d32a25b5344c72c9180ee7eaf6fda..e7c1f48dfd117508c5065116e281c53579b85319 100644 (file)
@@ -1,5 +1,5 @@
 CDIR=../../../..
-CHARMC=$(CDIR)/bin/charmc $(OPTS)
+CHARMC=$(CDIR)/bin/charmc $(OPTS) -I../../..
 FLAGS=-DEXTERIOR_BLACK_PIXEL_ELIMINATION
 
 HEADERS=liveViz.h liveViz.decl.h liveVizPoll.decl.h liveViz0.h colorScale.h ImageData.h