More changes.
authorIsaac Dooley <idooley2@illinois.edu>
Mon, 5 Mar 2007 21:51:17 +0000 (21:51 +0000)
committerIsaac Dooley <idooley2@illinois.edu>
Mon, 5 Mar 2007 21:51:17 +0000 (21:51 +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 [new file with mode: 0644]
src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS_CUDA.h [new file with mode: 0644]

index 705edf4d4253fb852451287ec8fbdf684df4f2de..deac18c94b7881cfdb909cb1dcdcd88d33ee9617 100644 (file)
 #include "ParFUM.decl.h"
 #include "ParFUM_internals.h"
 
-int elem_attr_size, node_attr_size;
-
 TopModel* topModel_Create_Init(int elem_attr_sz, int node_attr_sz){
   CkAssert(elem_attr_sz > 0);
   CkAssert(node_attr_sz > 0);
-  elem_attr_size = elem_attr_sz;
-  node_attr_size = node_attr_sz;
   TopModel *model = new TopModel;
+  model->elem_attr_size = elem_attr_sz;
+  model->node_attr_size = node_attr_sz;
 
   // This only uses a single mesh, so better not create multiple ones of these
   int which_mesh=FEM_Mesh_default_write();
@@ -37,14 +35,14 @@ TopModel* topModel_Create_Init(int elem_attr_sz, int node_attr_sz){
   // Allocate element connectivity
   FEM_Mesh_data(which_mesh,FEM_ELEM+0,FEM_CONN,temp_array, 0, 0, FEM_INDEX_0, 4);
   // Allocate element attributes
-  FEM_Mesh_data(which_mesh,FEM_ELEM+0,FEM_DATA+0,temp_array, 0, 0, FEM_BYTE, elem_attr_size);
+  FEM_Mesh_data(which_mesh,FEM_ELEM+0,FEM_DATA+0,temp_array, 0, 0, FEM_BYTE, model->elem_attr_size);
   // Allocate element Id array
   FEM_Mesh_data(which_mesh,FEM_ELEM+0,FEM_DATA+1,temp_array, 0, 0, FEM_INT, 1);
 
   // Allocate node coords
   FEM_Mesh_data(which_mesh,FEM_NODE,FEM_COORD,temp_array, 0, 0, FEM_DOUBLE, 3);
   // Allocate node attributes
-  FEM_Mesh_data(which_mesh,FEM_NODE+0,FEM_DATA+0,temp_array, 0, 0, FEM_BYTE, node_attr_size);
+  FEM_Mesh_data(which_mesh,FEM_NODE+0,FEM_DATA+0,temp_array, 0, 0, FEM_BYTE, model->node_attr_size);
   // Allocate node Id array
   FEM_Mesh_data(which_mesh,FEM_NODE+0,FEM_DATA+1,temp_array, 0, 0, FEM_INT, 1);
 
@@ -56,7 +54,6 @@ TopModel* topModel_Create_Init(int elem_attr_sz, int node_attr_sz){
   FEM_Mesh_allocate_valid_attr(which_mesh, FEM_NODE);
   FEM_Mesh_allocate_valid_attr(which_mesh, FEM_ELEM+0);
 
-
   // Setup ghost layers, assume Tet4
   // const int tet4vertices[4] = {0,1,2,3};
   //  FEM_Add_ghost_layer(1,1);
@@ -69,10 +66,11 @@ TopModel* topModel_Create_Driver(int elem_attr_sz, int node_attr_sz, int model_a
     // This only uses a single mesh, so don't create multiple TopModels of these
     CkAssert(elem_attr_sz > 0);
     CkAssert(node_attr_sz > 0);
-    elem_attr_size = elem_attr_sz;
-    node_attr_size = node_attr_sz;
     int which_mesh=FEM_Mesh_default_read();
     TopModel *model = new TopModel;
+    model->elem_attr_size = elem_attr_sz;
+    model->node_attr_size = node_attr_sz;
+    model->model_attr_size = model_attr_sz;
 
     model->mesh = FEM_Mesh_lookup(which_mesh,"TopModel::TopModel()");
 
@@ -99,9 +97,16 @@ TopModel* topModel_Create_Driver(int elem_attr_sz, int node_attr_sz, int model_a
 
     /** Copy model Attribute to device global memory */
     {
-        cudaMalloc(sizeof(ModelAtt), (void**)&(model->mAttDevice));
-        cudaMemcpy(mAttDevice,mAtt,sizeof(ModelAtt),cudaMemcpyHostToDevice);
+        cudaMalloc(model->model_attr_size, (void**)&(model->mAttDevice));
+        cudaMemcpy(mAttDevice,mAtt,model->model_attr_size,cudaMemcpyHostToDevice);
+    }
+
+    /** Copy model to device global memory */
+    {
+        cudaMalloc(sizeof(TopModel), (void**)&(model->modelD);
+        cudaMemcpy(model->modelD,model,sizeof(TopModel),cudaMemcpyHostToDevice);
     }
+
 #endif
 
     return model;
@@ -173,7 +178,7 @@ void topNode_SetAttrib(TopModel* m, TopNode n, void* d){
   FEM_DataAttribute * at = (FEM_DataAttribute*) m->mesh->node.lookup(FEM_DATA+0,"topNode_SetAttrib");
   AllocTable2d<unsigned char> &dataTable  = at->getChar();
   unsigned char *data = dataTable.getData();
-  memcpy(data + n*node_attr_size, d, node_attr_size);
+  memcpy(data + n*m->node_attr_size, d, m->node_attr_size);
 }
 
 /** @brief Set attribute of an element
@@ -184,7 +189,7 @@ void topElement_SetAttrib(TopModel* m, TopElement e, void* d){
   FEM_DataAttribute * at = (FEM_DataAttribute*) m->mesh->elem[0].lookup(FEM_DATA+0,"topElem_SetAttrib");
   AllocTable2d<unsigned char> &dataTable  = at->getChar();
   unsigned char *data = dataTable.getData();
-  memcpy(data + e*elem_attr_size, d, elem_attr_size);
+  memcpy(data + e*m->elem_attr_size, d, m->elem_attr_size);
 }
 
 
@@ -198,7 +203,7 @@ void* topElement_GetAttrib(TopModel* m, TopElement e){
   FEM_DataAttribute * at = (FEM_DataAttribute*) m->mesh->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 (data + e*m->elem_attr_size);
 }
 
 /** @brief Get nodal attribute
@@ -211,7 +216,7 @@ void* topNode_GetAttrib(TopModel* m, TopNode n){
   FEM_DataAttribute * at = (FEM_DataAttribute*) m->mesh->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 (data + n*m->node_attr_size);
 }
 
 
index fac7d72a9115bc1d59fa0cf29e73962596cbc29d..a54543c22b101d98853ce02c1f8f90ed6cb77875 100644 (file)
@@ -61,10 +61,15 @@ Sample usage:
 typedef struct{
     FEM_Mesh *mesh;
 
+    unsigned node_attr_size;
+    unsigned elem_attr_size;
+    unsigned model_attr_size;
+
 #ifdef CUDA
     unsigned char *mAttDevice; /** Device pointers to the goods */
     unsigned char *ElemDataDevice;
     unsigned char *NodeDataDevice;
+    TopModel* modelD;
 #endif
 
 } TopModel;
diff --git a/src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS_CUDA.C b/src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS_CUDA.C
new file mode 100644 (file)
index 0000000..9d2c7cf
--- /dev/null
@@ -0,0 +1,48 @@
+/**
+
+   @file
+   @brief Implementation of ParFUM-TOPS layer, except for Iterators
+
+   @author Isaac Dooley, Aaron Becker
+
+*/
+
+#include "ParFUM_TOPS.h"
+#include "ParFUM.decl.h"
+#include "ParFUM_internals.h"
+
+
+__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);
+}
+
+
+__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);
+}
+
+
+__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"
diff --git a/src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS_CUDA.h b/src/libs/ck-libs/ParFUM-Tops/ParFUM_TOPS_CUDA.h
new file mode 100644 (file)
index 0000000..e1c7969
--- /dev/null
@@ -0,0 +1,25 @@
+/**
+   @file
+   @brief A ParFUM "Tops" compatibility layer API Definition
+
+   @author Isaac Dooley and Aaron
+
+   ParFUM-TOPS provides a Tops-like API for ParFUM, meant to run on CUDA NVIDIA system.
+
+*/
+
+#ifndef __PARFUM_TOPS_CUDA___H
+#define __PARFUM_TOPS_CUDA___H
+
+#include <ParFUM.h>
+#include <ParFUM_internals.h>
+#include "ParFUM_TOPS.h"
+
+
+void* topElement_D_GetAttrib(TopModel* m, TopElement e);
+
+void* topNode_D_GetAttrib(TopModel* m, TopNode n);
+
+TopNode topElement_D_GetNode(TopModel* m,TopElement e,int idx);
+
+#endif