Go to the documentation of this file.00001 #ifndef LIBGEODECOMP_STORAGE_CUDAGRID_H
00002 #define LIBGEODECOMP_STORAGE_CUDAGRID_H
00003
00004 #include <libgeodecomp/geometry/coordbox.h>
00005 #include <libgeodecomp/geometry/region.h>
00006 #include <libgeodecomp/storage/cudaarray.h>
00007
00008 #include <cuda.h>
00009
00010 namespace LibGeoDecomp {
00011
00018 template<typename CELL_TYPE,
00019 typename TOPOLOGY=Topologies::Cube<2>::Topology,
00020 bool TOPOLOGICALLY_CORRECT=false>
00021 class CUDAGrid
00022 {
00023 public:
00024 friend class CUDAGridTest;
00025
00026 typedef CELL_TYPE CellType;
00027 typedef TOPOLOGY Topology;
00028
00029 static const int DIM = Topology::DIM;
00030
00031 inline CUDAGrid(
00032 const CoordBox<DIM>& box = CoordBox<DIM>(),
00033 const Coord<DIM>& topologicalDimensions = Coord<DIM>()) :
00034 box(box),
00035 topoDimensions(topologicalDimensions),
00036 array(box.dimensions.prod()),
00037 edgeCellStore(1)
00038 {}
00039
00040 void setEdge(const CELL_TYPE& cell)
00041 {
00042 cudaMemcpy(edgeCellStore.data(), &cell, sizeof(CELL_TYPE), cudaMemcpyHostToDevice);
00043 }
00044
00045 CELL_TYPE getEdge() const
00046 {
00047 CELL_TYPE cell;
00048 cudaMemcpy(&cell, edgeCellStore.data(), sizeof(CELL_TYPE), cudaMemcpyDeviceToHost);
00049 return cell;
00050 }
00051
00052 template<typename GRID_TYPE, typename REGION>
00053 void saveRegion(GRID_TYPE *target, const REGION& region) const
00054 {
00055 for (typename REGION::StreakIterator i = region.beginStreak();
00056 i != region.endStreak();
00057 ++i) {
00058 std::size_t length = i->length() * sizeof(CellType);
00059 cudaMemcpy(&(*target)[i->origin], const_cast<CellType*>(address(i->origin)), length, cudaMemcpyDeviceToHost);
00060 }
00061 }
00062
00063 template<typename GRID_TYPE, typename REGION>
00064 void loadRegion(const GRID_TYPE& source, const REGION& region)
00065 {
00066 for (typename REGION::StreakIterator i = region.beginStreak();
00067 i != region.endStreak();
00068 ++i) {
00069 std::size_t length = i->length() * sizeof(CellType);
00070 cudaMemcpy(address(i->origin), &source[i->origin], length, cudaMemcpyHostToDevice);
00071 }
00072 }
00073
00074 __host__ __device__
00075 const CoordBox<DIM>& boundingBox()
00076 {
00077 return box;
00078 }
00079
00080 __host__ __device__
00081 CellType *data()
00082 {
00083 return array.data();
00084 }
00085
00086 __host__ __device__
00087 const CellType *data() const
00088 {
00089 return array.data();
00090 }
00091
00092 __host__ __device__
00093 CellType *edgeCell()
00094 {
00095 return edgeCellStore.data();
00096 }
00097
00098 __host__ __device__
00099 const CellType *edgeCell() const
00100 {
00101 return edgeCellStore.data();
00102 }
00103
00104 private:
00105 CoordBox<DIM> box;
00106 Coord<DIM> topoDimensions;
00107 CUDAArray<CellType> array;
00108 CUDAArray<CellType> edgeCellStore;
00109
00110 std::size_t byteSize() const
00111 {
00112 return box.dimensions.prod() * sizeof(CellType);
00113 }
00114
00115 std::size_t offset(const Coord<DIM>& absoluteCoord) const
00116 {
00117 Coord<DIM> relativeCoord = absoluteCoord - box.origin;
00118 if (TOPOLOGICALLY_CORRECT) {
00119 relativeCoord = Topology::normalize(relativeCoord, topoDimensions);
00120 }
00121
00122 return relativeCoord.toIndex(box.dimensions);
00123 }
00124
00125 CellType *address(const Coord<DIM>& absoluteCoord)
00126 {
00127 return data() + offset(absoluteCoord);
00128 }
00129
00130 const CellType *address(const Coord<DIM>& absoluteCoord) const
00131 {
00132 return data() + offset(absoluteCoord);
00133 }
00134 };
00135
00136 }
00137
00138 #endif
