Changeset 162:e4e7373b7c14

Timestamp:

03/21/2012 01:47:32 PM (14 months ago)

Author:

Andreas Schaefer <gentryx@…>

Branch:

default

Message:

• refactored CanvasCell? interface so that it works with both, CPU and GPU simulators,
• GPUSimulator works, for now

Location:

src/examples/flowingcanvas

Files:

• CMakeLists.txt (modified) (1 diff)
• canvascell.h (modified) (6 diffs)
• canvasinitializer.h (modified) (1 diff)
• canvaswriter.h (modified) (1 diff)
• interactivesimulatorgpu.h (modified) (6 diffs)
• main.cu (moved) (moved from src/examples/flowingcanvas/main.cpp) (1 diff)

src/examples/flowingcanvas/CMakeLists.txt

@@ -21 +21 @@
 include_directories(./)
 
-add_executable(flowingcanvas ${SOURCES} ${MY_MOC_SOURCES})
-target_link_libraries(flowingcanvas ${LOCAL_LIBGEODECOMP_LINK_LIB})
+set(CUDA_TOOLKIT_ROOT_DIR "/opt/cuda")
+find_package(CUDA)
+find_package(Qt4)
+
+if (CUDA_FOUND)
+  cuda_add_executable(flowingcanvas ${SOURCES} ${MY_MOC_SOURCES})
+  target_link_libraries(flowingcanvas ${LOCAL_LIBGEODECOMP_LINK_LIB})
+endif(CUDA_FOUND)
+

src/examples/flowingcanvas/canvascell.h

@@ -4 +4 @@
 #include <libgeodecomp/misc/floatcoord.h>
 #include <libgeodecomp/misc/topologies.h>
+
+#ifndef __host__
+#define __host__
+#endif
+
+#ifndef __device__
+#define __device__
+#endif
 
 namespace LibGeoDecomp {
@@ -61 +69 @@
     }
 
-    template<typename COORD_MAP>
-    void update(const COORD_MAP& hood, const unsigned& nanoStep)
-    {
-        const CanvasCell& oldSelf = hood[Coord<2>()];
+    // fixme: faster conversion if moore-neighborhood is used (instead of von neumann)?
+    __host__ __device__
+    void update(const CanvasCell *up, const CanvasCell *same, const CanvasCell *down, const unsigned& nanoStep)
+    {
+        const CanvasCell& oldSelf = *same;
 
         pos[0] = oldSelf.pos[0];
@@ -78 +87 @@
         // fixme: move particles to other cells
         // fixme: kill dead particles
-        if (numParticles < 1) {
-            particles[numParticles] = Particle(pos[0], pos[1]);
-            numParticles = 1;
-        }
+        // if (numParticles < 1) {
+            // particles[numParticles] = Particle(pos[0], pos[1]);
+            // numParticles = 1;
+        // }
 
         forceSet = oldSelf.forceSet;
@@ -88 +97 @@
             forceFixed[1] = oldSelf.forceFixed[1];
         } else {
-            forceFixed[0] = (hood[Coord<2>(0, -1)].forceFixed[0] +
-                             hood[Coord<2>(-1, 0)].forceFixed[0] +
-                             hood[Coord<2>(1,  0)].forceFixed[0] +
-                             hood[Coord<2>(0,  1)].forceFixed[0]) * 0.25;
-            forceFixed[1] = (hood[Coord<2>(0, -1)].forceFixed[1] +
-                             hood[Coord<2>(-1, 0)].forceFixed[1] +
-                             hood[Coord<2>(1,  0)].forceFixed[1] +
-                             hood[Coord<2>(0,  1)].forceFixed[1]) * 0.25;
-        }
-
-        cameraLevel = (hood[Coord<2>(0, -1)].cameraLevel +
-                             hood[Coord<2>(-1, 0)].cameraLevel +
-                             hood[Coord<2>(1,  0)].cameraLevel +
-                             hood[Coord<2>(0,  1)].cameraLevel) * 0.25;
-        
-        float gradientX = hood[Coord<2>(1, 0)].cameraLevel - hood[Coord<2>(-1, 0)].cameraLevel;
-        float gradientY = hood[Coord<2>(0, 1)].cameraLevel - hood[Coord<2>(0, -1)].cameraLevel;
+            forceFixed[0] = (up[0].forceFixed[0] +
+                             same[-1].forceFixed[0] +
+                             same[1].forceFixed[0] +
+                             down[0].forceFixed[0]) * 0.25;
+            forceFixed[1] = (up[0].forceFixed[1] +
+                             same[-1].forceFixed[1] +
+                             same[1].forceFixed[1] +
+                             down[0].forceFixed[1]) * 0.25;
+        }
+
+        cameraLevel = (up[0].cameraLevel +
+                             same[-1].cameraLevel +
+                             same[1].cameraLevel +
+                             down[0].cameraLevel) * 0.25;
+        
+        float gradientX = same[1].cameraLevel - same[-1].cameraLevel;
+        float gradientY = down[0].cameraLevel - up[0].cameraLevel;
         forceVario[0] = 0;
         if ((gradientX > 0.011) || (gradientX < -0.011)) {
@@ -122 +131 @@
 
         for (int i = 0; i < numParticles; ++i) {
-            Particle& p = particles[i];
+            // Particle& p = particles[i];
             // fixme: parameters
-            p.update(1.0, forceTotal[0], forceTotal[1], 1.0, 0.99);
+            // p.update(1.0, forceTotal[0], forceTotal[1], 1.0, 0.99);
         }
         
@@ -146 +155 @@
 //                         forceVario[1] + forceFixed[1]);
 //         // moveParticles();
+    }
+
+    template<typename COORD_MAP>
+    void update(const COORD_MAP& hood, const unsigned& nanoStep)
+    {
+        update(&hood[Coord<2>(0, -1)], &hood[Coord<2>(0, 0)], &hood[Coord<2>(0, 1)], nanoStep);
     }
 

src/examples/flowingcanvas/canvasinitializer.h

@@ -12 +12 @@
     CanvasInitializer() :
         // SimpleInitializer<CanvasCell>(Coord<2>(240, 135), 100)
-        SimpleInitializer<CanvasCell>(Coord<2>(320, 180), 100)
+        // SimpleInitializer<CanvasCell>(Coord<2>(320, 180), 100)
+        SimpleInitializer<CanvasCell>(Coord<2>(384, 216), 100)
         // SimpleInitializer<CanvasCell>(Coord<2>(640, 360), 100)
     {}

src/examples/flowingcanvas/canvaswriter.h

@@ -61 +61 @@
     CanvasWriter(QImage **_outputFrame,
                  MonolithicSimulator<CanvasCell> *_sim) :
-        Writer("foo", _sim, 1),
+        Writer<CanvasCell>("foo", _sim, 1),
         outputFrame(_outputFrame),
-        mode(5)
+        mode(4)
     {}
 

src/examples/flowingcanvas/interactivesimulatorgpu.h

@@ -5 +5 @@
 #include <libgeodecomp/misc/grid.h>
 
+__global__ void updateKernel(CanvasCell *curGrid, CanvasCell *newGrid, unsigned width, unsigned nanoStep)
+{
+    int x = 1 + blockDim.x * blockIdx.x + threadIdx.x;
+    int y = 1 + blockDim.y * blockIdx.y + threadIdx.y;
+    int index = y * width + x;
+
+    newGrid[index].update(
+        curGrid + index - width, 
+        curGrid + index,
+        curGrid + index + width,
+        nanoStep);
+}
+
+__global__ void loadGridFromTransferBuffer(CanvasCell *grid, CanvasCell *buffer, unsigned widthGrid, unsigned widthBuffer)
+{
+    int x = blockDim.x * blockIdx.x + threadIdx.x;
+    int y = blockDim.y * blockIdx.y + threadIdx.y;
+    int indexGrid   = (y + 1) * widthGrid + x + 1;
+    int indexBuffer = y * widthBuffer + x;
+
+    grid[indexGrid] = buffer[indexBuffer];
+}
+
+__global__ void storeGridToTransferBuffer(CanvasCell *grid, CanvasCell *buffer, unsigned widthGrid, unsigned widthBuffer)
+{
+    int x = blockDim.x * blockIdx.x + threadIdx.x;
+    int y = blockDim.y * blockIdx.y + threadIdx.y;
+    int indexGrid   = (y + 1) * widthGrid + x + 1;
+    int indexBuffer = y * widthBuffer + x;
+
+    buffer[indexBuffer] = grid[indexGrid];
+}
+
 namespace LibGeoDecomp {
 
+// fixme: this is hardcoded to 2d ATM
 template<typename CELL_TYPE>
 class GPUSimulator : public MonolithicSimulator<CELL_TYPE>
@@ -13 +47 @@
     typedef typename CELL_TYPE::Topology Topology;
     typedef Grid<CELL_TYPE, Topology> GridType;
-    static const int DIMENSIONS = Topology::DIMENSIONS;
-
-    GPUSimulator(Initializer<CELL_TYPE> *initializer) :
+    static const int DIM = Topology::DIMENSIONS;
+
+    GPUSimulator(Initializer<CELL_TYPE> *initializer, const int& device = 0) :
         MonolithicSimulator<CELL_TYPE>(initializer)
     {
-        gridHost.resize(this->initializer->gridBox().dimensions);
+        Coord<DIM> dim = this->initializer->gridDimensions();
+        gridHost.resize(dim);
         this->initializer->grid(&gridHost);
+
+        cudaSetDevice(device);
+        int byteSize = dim.prod() * sizeof(CELL_TYPE);
+        cudaMalloc(&transferGrid, byteSize);
+        cudaMemcpy(transferGrid, gridHost.baseAddress(), byteSize, cudaMemcpyHostToDevice);
+
+        // pad actual grids to avoid edge cell handling
+        Coord<2> paddedDim = dim + Coord<2>(2, 2);
+        byteSize = paddedDim.prod() * sizeof(CELL_TYPE);
+        cudaMalloc(&curGridDevice, byteSize);
+        cudaMalloc(&newGridDevice, byteSize);
+        GridType initGrid(paddedDim, gridHost.getEdgeCell(), gridHost.getEdgeCell());
+        cudaMemcpy(curGridDevice, initGrid.baseAddress(), byteSize, cudaMemcpyHostToDevice);
+        cudaMemcpy(newGridDevice, initGrid.baseAddress(), byteSize, cudaMemcpyHostToDevice);
+
+        dim3 blockDim;
+        dim3 gridDim;
+        genKernelDimensions(&blockDim, &gridDim);
+        loadGridFromTransferBuffer<<<gridDim, blockDim>>>(curGridDevice, transferGrid, gridWidth() + 2, gridWidth());
+        checkCudaError();
+    }
+
+    virtual ~GPUSimulator()
+    {
+        cudaFree(&curGridDevice);
+        cudaFree(&newGridDevice);
     }
 
@@ -40 +101 @@
     virtual const GridType *getGrid()
     {
+        dim3 blockDim;
+        dim3 gridDim;
+        genKernelDimensions(&blockDim, &gridDim);
+        storeGridToTransferBuffer<<<gridDim, blockDim>>>(curGridDevice, transferGrid, gridWidth() + 2, gridWidth());
+        int byteSize = gridHost.getDimensions().prod() * sizeof(CELL_TYPE);
+        cudaMemcpy(gridHost.baseAddress(), transferGrid, byteSize, cudaMemcpyDeviceToHost);
+        checkCudaError();
         return &gridHost;
     }
@@ -45 +113 @@
 private:
     GridType gridHost;
+    CELL_TYPE *transferGrid;
     CELL_TYPE *curGridDevice;
     CELL_TYPE *newGridDevice;
 
+    void checkCudaError()
+    {
+        cudaError_t error = cudaGetLastError();
+        if (error != cudaSuccess) {
+            const char *errorMessage = cudaGetErrorString(error);
+            std::cerr << "CUDA: " << errorMessage << "\n";
+            throw std::runtime_error("CUDA call failed");
+        }
+    }
+
     void nanoStep(const unsigned& nanoStep)
     {
-        std::cout << "fixme InteractiveSimulatorGPU::nanoStep()\n";
+        dim3 blockDim;
+        dim3 gridDim;
+        genKernelDimensions(&blockDim, &gridDim);
+        updateKernel<<<gridDim, blockDim>>>(curGridDevice, newGridDevice, gridWidth() + 2, nanoStep);
+        checkCudaError();
+        std::swap(curGridDevice, newGridDevice);
+    }
+
+    void genKernelDimensions(dim3 *blockDim, dim3 *gridDim)
+    {
+        Coord<DIM> dim = gridHost.getDimensions();
+        int blockDimX = 32;
+        int blockDimY = 8;
+        *blockDim = dim3(blockDimX, blockDimY);
+        *gridDim = dim3(dim.x() / blockDimX, dim.y() / blockDimY);
+    }
+
+    int gridWidth()
+    {
+        return gridHost.getDimensions().x();
     }
 };
@@ -60 +158 @@
     typedef typename CELL_TYPE::Topology Topology;
     typedef Grid<CELL_TYPE, Topology> GridType;
-    static const int DIMENSIONS = Topology::DIMENSIONS;
+    typedef std::vector<boost::shared_ptr<Writer<CELL_TYPE> > > WriterVector;
 
     InteractiveSimulatorGPU(QObject *parent, Initializer<CELL_TYPE> *initializer) :
@@ -77 +175 @@
     virtual void renderOutput()
     {
-        std::cout << "fixme InteractiveSimulatorGPU::renderOutput()\n";
+        // fixme: this is the same for InteractiveSimulatorCPU. refactor?
+        for(unsigned i = 0; i < this->writers.size(); i++) 
+            this->writers[i]->stepFinished();
     }
 
     virtual void update()
     {
+        // for (int fixme = 0; fixme < 500; ++fixme) 
         GPUSimulator<CELL_TYPE>::step();
-    }
+        // sleep(1);
+    }
+
+    virtual void registerWriter(Writer<CELL_TYPE> *writer)
+    {
+        writers.push_back(boost::shared_ptr<Writer<CELL_TYPE> >(writer));
+    }
+
+protected:
+    WriterVector writers;
 };
 

src/examples/flowingcanvas/main.cu

@@ -27 +27 @@
     flow.resize(1200, 900);
 
-    InteractiveSimulatorGPU<CanvasCell> *sim = new InteractiveSimulatorGPU<CanvasCell>(
+    // InteractiveSimulatorGPU<CanvasCell> *sim = new InteractiveSimulatorGPU<CanvasCell>(
+    InteractiveSimulatorCPU<CanvasCell> *sim = new InteractiveSimulatorCPU<CanvasCell>(
         &flow,
         new CanvasInitializer());