Some changes to make RR or Block mapping of PE to node work.
authorIsaac Dooley <idooley2@illinois.edu>
Tue, 7 Apr 2009 01:40:53 +0000 (01:40 +0000)
committerIsaac Dooley <idooley2@illinois.edu>
Tue, 7 Apr 2009 01:40:53 +0000 (01:40 +0000)
src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS.cc
src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS.h
src/libs/ck-libs/ParFUM/ParFUM_internals.h
src/libs/ck-libs/ParFUM/import.C
src/libs/ck-libs/ParFUM/import.h
src/libs/ck-libs/ParFUM/mesh.C

index aeef191ea4c845de735b74dba9003d0b70a77a19..d71324cfafd660855c154946455ed7eefff57e50 100644 (file)
@@ -309,7 +309,9 @@ TopModel* topModel_Create_Driver(TopDevice target_device, int elem_attr_sz,
     if (model->target_device == DeviceGPU) {
         int size = model->num_local_elem * connSize *sizeof(int);
         cudaError_t err = cudaMalloc((void**)&(model->device_model.n2eConnDevice), size);
-       if(err != cudaSuccess){
+       if(err == cudaErrorMemoryAllocation){
+                 CkPrintf("[%d] cudaMalloc FAILED with error cudaErrorMemoryAllocation model->device_model.n2eConnDevice in ParFUM_TOPS.cc size=%d: %s\n", CkMyPe(), size, cudaGetErrorString(err));
+       }else if(err != cudaSuccess){
          CkPrintf("[%d] cudaMalloc FAILED model->device_model.n2eConnDevice in ParFUM_TOPS.cc size=%d: %s\n", CkMyPe(), size, cudaGetErrorString(err));
          CkAbort("cudaMalloc FAILED");
        }
@@ -1097,6 +1099,7 @@ public:
   
   char *locations;
   int objs_per_block;
+  int numNodes;
 
   /// labels for states used when parsing the ConfigurableRRMap from ARGV
   enum loadStatus{
@@ -1137,12 +1140,10 @@ public:
        std::istringstream instream(configuration);
        CkAssert(instream.good());
         
-       // Example line:
-       // 5 C G G G C
-       // Modulo 5; The first and fifth VPs are on CPU, the second, third, and forth are on GPUs. 
 
        // extract first integer
        instream >> objs_per_block;
+       instream >> numNodes;
        CkAssert(instream.good());
        CkAssert(objs_per_block > 0);
        locations = new char[objs_per_block];
@@ -1181,6 +1182,11 @@ bool haveConfigurableCPUGPUMap(){
   return loader.haveConfiguration();
 }
 
+int configurableCPUGPUMapNumNodes(){
+  ConfigurableCPUGPUMapLoader &loader =  CkpvAccess(myConfigGPUCPUMapLoader);
+  return loader.numNodes;
+}
+
 
 bool isPartitionCPU(int partition){
   ConfigurableCPUGPUMapLoader &loader =  CkpvAccess(myConfigGPUCPUMapLoader);
index a2d004828c5fbd651cea7255b6951dc18392abce..70b610e895273c220aad8bc6ba837d84dde711c8 100644 (file)
@@ -274,7 +274,7 @@ TopElement topModel_InsertCohesiveAtFacet (TopModel* m, int ElemType, TopFacet f
 bool haveConfigurableCPUGPUMap();
 bool isPartitionCPU(int partition);
 bool isPartitionGPU(int partition);
-
+int configurableCPUGPUMapNumNodes();
 
 
 
index 6862fd94a4a6b2b3fadbea66d694281e5c5a2038..b953442cfbf3d06236834181f22d5ebf6a263460 100644 (file)
@@ -839,6 +839,9 @@ class FEM_Entity {
    */
   void setLength(int newlen, bool f=false);
 
+  /// Attempt to set max and resize all associated attributes
+  void reserveLength(int newlen);
+
   /** Support for registration API
    *  Set the current length and maximum length for this entity.
    *  If the current length exceeds the maximum length a resize
index b57f48f845e32d7ba02668f52ba6bf0d30b24a23..b9c07a52b641522cd01521b0a01ae78a18c3cbca 100644 (file)
@@ -18,6 +18,7 @@ void ParFUM_deghosting(int meshid){
 
 /// Recreate the shared nodes. An alternate incorrect version can be enabled by undefining CORRECT_COORD_COMPARISON
 void ParFUM_recreateSharedNodes(int meshid, int dim, MPI_Comm newComm) {
+
 #define CORRECT_COORD_COMPARISON
   MPI_Comm comm = newComm;
   int rank, nParts;
@@ -25,6 +26,43 @@ void ParFUM_recreateSharedNodes(int meshid, int dim, MPI_Comm newComm) {
   int recv_count=0; // sanity check
   MPI_Comm_rank(comm, &rank);
   MPI_Comm_size(comm, &nParts);
+
+
+#if SUPER_FAST_SPECIFIC_TORUS
+
+#define TORUSY 15
+#define TORUSZ 15
+
+  CkPrintf("rank %d is manually configuring the IDXL lists to make the shared node generation fast\n");
+  
+  FEM_Mesh *mesh = (FEM_chunk::get("ParFUM_recreateSharedNodes"))->lookup(meshid,"ParFUM_recreateSharedNodes");
+  IDXL_Side &shared = mesh->node.shared;
+  
+  int low = (rank-1+nParts) % nParts;
+  int high = (rank+1) % nParts;
+  
+  IDXL_List &list1 = shared.addList(low);
+  IDXL_List &list2 = shared.addList(high);
+  
+  int nodesInPlane = TORUSY * TORUSZ;
+  int numNodes = FEM_Mesh_get_length(meshid,FEM_NODE);
+  
+  // vp - 1
+  for(int j=0;j<nodesInPlane;j++){
+    list1.push_back(j);
+  }
+  
+  // vp + 1
+  for(int j=0;j<nodesInPlane;j++){
+    list2.push_back(numNodes - nodesInPlane +j);
+  }
+  
+  return;
+#else
+
+
+
+
   // Shared data will be temporarily stored in the following structure
   int *sharedNodeCounts; // sharedCounts[i] = number of nodes shared with rank i
   int **sharedNodeLists; // sharedNodes[i] is the list of nodes shared with rank i
@@ -102,6 +140,86 @@ void ParFUM_recreateSharedNodes(int meshid, int dim, MPI_Comm newComm) {
     // Match coords between local nodes and received coords
     int recvNodeCount = length/dim;
 
+
+    // PERFORM THE NODE COMPARISONS
+
+#ifdef SHARED_NODES_ONLY_NEIGHBOR
+
+    int borderNodes = BORDERNODES;
+
+#warning "Only the first and last BORDERNODES nodes on each partition are candidates for being shared nodes"
+
+    // indices are inclusive
+    int myBottomLow = 0;
+    int myBottomHigh = borderNodes;
+    int myTopLow = numNodes - borderNodes;
+    int myTopHigh = numNodes-1;
+
+    int recvBottomLow = 0;
+    int recvBottomHigh = borderNodes;
+    int recvTopLow = recvNodeCount - borderNodes;
+    int recvTopHigh = recvNodeCount-1;
+
+    CkPrintf("[%d] rank=%d myBottomLow=%d myBottomHigh=%d myTopLow=%d myTopHigh=%d   recvBottomLow=%d recvBottomHigh=%d recvTopLow=%d recvTopHigh=%d\n", CkMyPe(), rank, myBottomLow, myBottomHigh, myTopLow, myTopHigh, recvBottomLow, recvBottomHigh, recvTopLow, recvTopHigh);    
+
+    // make sure the top region is non-negative
+    if(myTopLow < 0)
+      myTopLow = 0;
+      
+    if(recvTopLow < 0)
+      recvTopLow = 0;
+
+    // make the two regions be non-overlapping
+    if(myBottomHigh >= myTopLow)
+      myTopLow = myTopLow-1;
+    
+    if(recvBottomHigh >= recvTopLow)
+      recvTopLow = recvTopLow-1;
+    
+    for (int j=myBottomLow; j<=myBottomHigh; j++) {
+      for (int k=recvBottomLow; k<=recvBottomHigh; k++) {
+       if (coordEqual(&nodeCoords[j*dim], &recvNodeCoords[k*dim], dim)) {
+         localSharedNodes.push_back(j); 
+         remoteSharedNodes.push_back(k);
+         break;
+       }
+      }
+    }
+
+    for (int j=myTopLow; j<=myBottomHigh; j++) {
+      for (int k=recvTopLow; k<=recvTopHigh; k++) {
+       if (coordEqual(&nodeCoords[j*dim], &recvNodeCoords[k*dim], dim)) {
+         localSharedNodes.push_back(j); 
+         remoteSharedNodes.push_back(k);
+         break;
+       }
+      }
+    }
+
+
+    for (int j=myTopLow; j<=myTopHigh; j++) {
+      for (int k=recvBottomLow; k<=recvBottomHigh; k++) {
+       if (coordEqual(&nodeCoords[j*dim], &recvNodeCoords[k*dim], dim)) {
+         localSharedNodes.push_back(j); 
+         remoteSharedNodes.push_back(k);
+         break;
+       }
+      }
+    }
+
+    for (int j=myBottomLow; j<=myTopHigh; j++) {
+      for (int k=recvTopLow; k<=recvTopHigh; k++) {
+       if (coordEqual(&nodeCoords[j*dim], &recvNodeCoords[k*dim], dim)) {
+         localSharedNodes.push_back(j); 
+         remoteSharedNodes.push_back(k);
+         break;
+       }
+      }
+    }
+
+#else 
+
+
     //    CkPrintf("Comparing %d nodes with %d received nodes\n", numNodes, recvNodeCount);
     for (int j=0; j<numNodes; j++) {
       for (int k=0; k<recvNodeCount; k++) {
@@ -114,6 +232,8 @@ void ParFUM_recreateSharedNodes(int meshid, int dim, MPI_Comm newComm) {
       }
     }
 
+#endif
+
     // Copy local nodes that were shared with source into the data structure
     int *localSharedNodeList = (int *)malloc(localSharedNodes.size()*sizeof(int));
     for (int m=0; m<localSharedNodes.size(); m++) {
@@ -185,6 +305,8 @@ void ParFUM_recreateSharedNodes(int meshid, int dim, MPI_Comm newComm) {
       free(sharedNodeLists[i]);
   }
   free(sharedNodeLists);
+
+#endif // normal mode, not super fast mesh specific one
 }
 
 void ParFUM_createComm(int meshid, int dim, MPI_Comm comm)
index bad4d5a71a919cbf3537489290b8b94b9a084e50..14794464e9cb1209f0ff7d4fdaa095c9c64bffbf 100644 (file)
@@ -42,7 +42,7 @@ inline int coordCompare(const double *key1, const double *key2, int dim) {
   return 0;
 #else
   CkAssert(dim==3);
-  const double threshold = 0.001;
+  const double threshold = 0.0001;
   for(int i=0; i<dim; i++) {
     const double a = key1[i];
     const double b = key2[i];
index 8c5e6f87d02e497290292c54a8f337f074ca218c..5d19386ab18bbd812c492808600edf8f293dfb3e 100644 (file)
@@ -1220,6 +1220,19 @@ void FEM_Entity::setLength(int newlen, bool f)
     }
 }
 
+/// Attempt to resize all the attributes owned by ParFUM so they don't need to be resized later. 
+/// This function is not guaranteed to resize all buffers. The usual resize methods should be used 
+/// later on, but they should be cheaper at that point.
+void FEM_Entity::reserveLength(int newlen)
+{
+  if (newlen > max) {
+    max = newlen;
+    for (int a=0;a<attributes.size();a++){
+      attributes[a]->reallocate();
+    }
+  }
+}
+
 void FEM_Entity::allocateValid(void) {
   if (!valid){
     valid=new FEM_DataAttribute(this,FEM_IS_VALID_ATTR);