More changes!
authorIsaac Dooley <idooley2@illinois.edu>
Mon, 5 Mar 2007 22:56:53 +0000 (22:56 +0000)
committerIsaac Dooley <idooley2@illinois.edu>
Mon, 5 Mar 2007 22:56:53 +0000 (22:56 +0000)
src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS.C
src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS.h
src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS_CUDA.C

index deac18c94b7881cfdb909cb1dcdcd88d33ee9617..ad9324d6615ffc14657b0df83e1626c66c8e73f5 100644 (file)
@@ -74,6 +74,8 @@ TopModel* topModel_Create_Driver(int elem_attr_sz, int node_attr_sz, int model_a
 
     model->mesh = FEM_Mesh_lookup(which_mesh,"TopModel::TopModel()");
 
+    model->mAtt = mAtt;
+
 #if CUDA
     /** Copy element Attribute array to device global memory */
     {
@@ -83,6 +85,7 @@ TopModel* topModel_Create_Driver(int elem_attr_sz, int node_attr_sz, int model_a
         int size = at->size();
         cudaMalloc(size, (void**)&(model->ElemDataDevice));
         cudaMemcpy(ElemDataDevice,ElemData,size,cudaMemcpyHostToDevice);
+        model->num_local_elem = size / elem_attr_sz;
     }
 
     /** Copy node Attribute array to device global memory */
@@ -93,6 +96,17 @@ TopModel* topModel_Create_Driver(int elem_attr_sz, int node_attr_sz, int model_a
         int size = at->size();
         cudaMalloc(size, (void**)&(model->NodeDataDevice));
         cudaMemcpy(NodeDataDevice,NodeData,size,cudaMemcpyHostToDevice);
+        model->num_local_node = size / node_attr_sz;
+    }
+
+    /** Copy elem connectivity array to device global memory */
+    {
+        FEM_DataAttribute * at = (FEM_DataAttribute*) m->mesh->node.lookup(FEM_CONN,"topModel_Create_Driver");
+        AllocTable2d<int> &dataTable  = at->getInt();
+        int *data = dataTable.getData();
+        int size = at->size()*sizeof(int);
+        cudaMalloc(size, (void**)&(model->ElemConnDevice));
+        cudaMemcpy(ElemConnDevice,data,size,cudaMemcpyHostToDevice);
     }
 
     /** Copy model Attribute to device global memory */
index a54543c22b101d98853ce02c1f8f90ed6cb77875..b7cd5831f285b0517a08c13a39501a1a2cb2ccc2 100644 (file)
@@ -60,18 +60,27 @@ Sample usage:
 /** A tops model is roughly equivalent to a ParFUM FEM_Mesh object */
 typedef struct{
     FEM_Mesh *mesh;
+    void *mAtt;
 
     unsigned node_attr_size;
     unsigned elem_attr_size;
     unsigned model_attr_size;
 
+    /** number of local elements */
+    unsigned num_local_elem;
+    /** number of local nodes */
+    unsigned num_local_node;
+
 #ifdef CUDA
     unsigned char *mAttDevice; /** Device pointers to the goods */
     unsigned char *ElemDataDevice;
     unsigned char *NodeDataDevice;
+    int * ElemConnDevice;
     TopModel* modelD;
 #endif
 
+
+
 } TopModel;
 
 
@@ -102,6 +111,9 @@ enum {
   TOP_ELEMENT_MAX
 };
 
+/** used as iterators on CUDA system. See usage!*/
+typedef bool TopNodeItr_D;
+typedef bool TopElemItr_D;
 
 /** Node Iterator */
 class TopNodeItr{
@@ -119,6 +131,7 @@ public:
     TopModel *model;
 };
 
+
 /** an opaque id for top entities */
 typedef int TopID;
 
index 9d2c7cff6038732fbf60d3f9aa141725c1ee2791..8451773348af6513d64a3f5e3049306186aac313 100644 (file)
 
 
 __device__ void* topElement_D_GetAttrib(TopModel* m, TopElement e){
-  if(! m->elem[0].is_valid_any_idx(e))
-       return NULL;
-
-  FEM_DataAttribute * at = (FEM_DataAttribute*) m->elem[0].lookup(FEM_DATA+0,"topElem_GetAttrib");
-  AllocTable2d<unsigned char> &dataTable  = at->getChar();
-  unsigned char *data = dataTable.getData();
-  return (data + e*elem_attr_size);
+  return (m->ElemDataDevice + e*m->elem_attr_size);
 }
 
 
 __device__ void* topNode_D_GetAttrib(TopModel* m, TopNode n){
-  if(! m->node.is_valid_any_idx(n))
-       return NULL;
-
-  FEM_DataAttribute * at = (FEM_DataAttribute*) m->node.lookup(FEM_DATA+0,"topNode_GetAttrib");
-  AllocTable2d<unsigned char> &dataTable  = at->getChar();
-  unsigned char *data = dataTable.getData();
-  return (data + n*node_attr_size);
+  return (m->NodeDataDevice + n*m->node_attr_size);
 }
 
 
-__device__ TopNode topElement_D_GetNode(TopModel* m,TopElement e,int idx){
-  CkAssert(e>=0);
-  const AllocTable2d<int> &conn = m->elem[0].getConn();
-  CkAssert(idx>=0 && idx<conn.width());
-  CkAssert(idx<conn.size());
-
-  int node = conn(e,idx);
-
-  return conn(e,idx);
-}
-
 #include "ParFUM_TOPS.def.h"