SWE
/import/home/rettenbs/src/SWE/src/blocks/cuda/SWE_BlockCUDA.hh
Go to the documentation of this file.
00001 
00029 #ifndef __SWE_BLOCKCUDA_HH
00030 #define __SWE_BLOCKCUDA_HH
00031 
00032 #include "blocks/SWE_Block.hh"
00033 
00034 #include "tools/help.hh"
00035 
00036 #include <iostream>
00037 #include <fstream>
00038 
00039 #include <cuda_runtime.h>
00040 
00041 using namespace std;
00042 
00043 void checkCUDAError(const char *msg);
00044 void tryCUDA(cudaError_t err, const char *msg);
00045 
00046 const int TILE_SIZE=16;
00047 //const int TILE_SIZE=8;
00048 
00055 class SWE_BlockCUDA : public SWE_Block {
00056 
00057   public:
00058     // Constructor und Destructor
00059     SWE_BlockCUDA(int l_nx, int l_ny,
00060                 float l_dx, float l_dy);
00061     virtual ~SWE_BlockCUDA();
00062     
00063   // object methods
00064 
00065 // ---> COULD BE IMPLEMENTED TO PROVIDE A DEFAULT IMPLEMENTATION
00066 //     // determine maximum possible time step
00067 //     virtual float getMaxTimestep();
00068 
00069     // deliver a pointer to proxy class that represents
00070     // the layer that is copied to an external ghost layer 
00071     virtual SWE_Block1D* registerCopyLayer(BoundaryEdge edge);
00072     // "grab" the ghost layer in order to set these values externally
00073     virtual SWE_Block1D* grabGhostLayer(BoundaryEdge edge);
00074 
00075     // access to CUDA variables
00079     const float* getCUDA_waterHeight() { return hd; };
00083     const float* getCUDA_bathymetry() { return bd; };
00084 
00085   protected:
00086      
00087     // synchronisation Methods
00088     virtual void synchAfterWrite();
00089     virtual void synchWaterHeightAfterWrite();
00090     virtual void synchDischargeAfterWrite();
00091     virtual void synchBathymetryAfterWrite();
00092     virtual void synchGhostLayerAfterWrite();
00093 
00094     virtual void synchBeforeRead();
00095     virtual void synchWaterHeightBeforeRead();
00096     virtual void synchDischargeBeforeRead();
00097     virtual void synchBathymetryBeforeRead();
00098     virtual void synchCopyLayerBeforeRead();
00099     
00100     // set boundary conditions in ghost layers (set boundary conditions)
00101     virtual void setBoundaryConditions();
00102 
00103     // define arrays for main unknowns in CUDA global memory: 
00104     // hd, hud, hvd, and bd are CUDA arrays corresp. to h, hu, hv, and b
00105     float* hd;
00106     float* hud;
00107     float* hvd;
00108     float* bd;
00109         
00110   private:
00111      
00112     // separate memory to hold bottom and top ghost and copy layer 
00113     // in main memory allowing non-strided access
00114     float* bottomLayer;
00115     float* topLayer;
00116     SWE_Block1D* bottomGhostLayer;
00117     SWE_Block1D* bottomCopyLayer;
00118     SWE_Block1D* topGhostLayer;
00119     SWE_Block1D* topCopyLayer;
00120     // and resp. memory on the CUDA device:
00121     float* bottomLayerDevice;
00122     float* topLayerDevice;
00123 
00124     // helper arrays: store maximum height and velocities to determine time step
00125     float* maxhd;
00126     float* maxvd;
00127 
00128   public:
00129     // print information about the CUDA device
00130     static void printDeviceInformation();
00131 
00136     static void init(int i_cudaDevice = 0);
00138     static void finalize();
00139 };
00140 
00146 inline __device__
00147 int getCellCoord(int x, int y, int ny) {
00148    return x*(ny+2) + y;
00149 }
00150 
00151 
00157 inline __device__
00158 int getEdgeCoord(int x, int y, int ny) {
00159    return x*(ny+1) + y;
00160 }
00161 
00167 inline __device__
00168 int getBathyCoord(int x, int y, int ny) {
00169    return x*ny + y;
00170 }
00171 
00172 
00173 #endif
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends