Changeset 163:1aae525e1e79

Timestamp:

03/21/2012 03:26:46 PM (15 months ago)

Author:

Andreas Schaefer <gentryx@…>

Branch:

default

Message:

camera frames are now propated to the GPU code

Location:

src/examples/flowingcanvas

Files:

• canvascell.h (modified) (3 diffs)
• interactivesimulatorcpu.h (modified) (1 diff)
• interactivesimulatorgpu.h (modified) (10 diffs)
• main.cu (modified) (1 diff)

src/examples/flowingcanvas/canvascell.h

@@ -21 +21 @@
     {
     public:
+        __host__ __device__
         Particle(const float& pos0 = 0, const float& pos1 = 0, const float& _lifetime = 1000) :
             lifetime(_lifetime)
@@ -30 +31 @@
         }
 
+        __host__ __device__
         void update(const float& deltaT, const float& force0, const float& force1, const float& forceFactor, const float& friction)
         {
@@ -163 +165 @@
     }
 
-    void readCam(unsigned char *frame, const float& factorX, const float& factorY, const int& width, const int& height)
-    {
-        int posX = pos[0] * factorX;
-        int posY = pos[1] * factorY;
-        
-        int index = posY * width * 3 + posX * 3;
-        cameraPixel = (0xff << 24) + 
-            (frame[index + 0] << 16) +
-            (frame[index + 1] <<  8) +
-            (frame[index + 2] <<  0);
-
-        int val = (int)frame[index + 0] + (int)frame[index + 1] + (int)frame[index + 2];
+
+    __host__ __device__
+    void readCam(const unsigned char& r, const unsigned char& g, const unsigned char& b)
+    {
+        cameraPixel = (0xff << 24) + (r << 16) + (g <<  8) + (b <<  0);
+
+        int val = (int)r + (int)g + (int)b;
         if (val < 500) {
             cameraLevel = 1.0;
         } else {
-            cameraLevel = std::max(0.0, cameraLevel - 0.05);
-        }
+            cameraLevel = max(0.0, cameraLevel - 0.05);
+        }
+    }
+
+    __host__ __device__
+    float max(const float& a, const float& b)
+    {
+        return a > b ? a : b;
     }
 

src/examples/flowingcanvas/interactivesimulatorcpu.h

@@ -33 +33 @@
             for (int x = 0; x < dim.x(); ++x) {
                 Coord<2> c(x, y);
-                (*this->curGrid)[c].readCam(&cameraFrame[0], factorX, factorY, cameraFrameWidth, cameraFrameHeight);
+                int index = (int)(y * factorY) * cameraFrameWidth + x * factorX;
+                unsigned char *pixel = &cameraFrame[3 * index];
+                (*this->curGrid)[c].readCam(pixel[0], pixel[1], pixel[2]);
             }
         }

src/examples/flowingcanvas/interactivesimulatorgpu.h

@@ -5 +5 @@
 #include <libgeodecomp/misc/grid.h>
 
+// fixme: kernels need to be templates. should be included by GPUSimulator
 __global__ void updateKernel(CanvasCell *curGrid, CanvasCell *newGrid, unsigned width, unsigned nanoStep)
 {
@@ -38 +39 @@
 }
 
+
+__global__ void updateCam(CanvasCell *grid, unsigned char *cameraBuffer, float factorX, float factorY, 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 = (int)(y * factorY) * widthBuffer + x * factorX;
+    unsigned char *pixel = cameraBuffer + indexBuffer * 3;
+   
+    grid[indexGrid].readCam(pixel[0], pixel[1], pixel[2]);
+}
+
 namespace LibGeoDecomp {
 
 // fixme: this is hardcoded to 2d ATM
+// fixme: move to dedicated file
 template<typename CELL_TYPE>
 class GPUSimulator : public MonolithicSimulator<CELL_TYPE>
@@ -50 +64 @@
 
     GPUSimulator(Initializer<CELL_TYPE> *initializer, const int& device = 0) :
-        MonolithicSimulator<CELL_TYPE>(initializer)
+        MonolithicSimulator<CELL_TYPE>(initializer),
+        inputBufferDevice(0),
+        inputBufferSize(0)
     {
         Coord<DIM> dim = this->initializer->gridDimensions();
@@ -58 +74 @@
         cudaSetDevice(device);
         int byteSize = dim.prod() * sizeof(CELL_TYPE);
-        cudaMalloc(&transferGrid, byteSize);
-        cudaMemcpy(transferGrid, gridHost.baseAddress(), byteSize, cudaMemcpyHostToDevice);
+        cudaMalloc(&transferGridDevice, byteSize);
+        cudaMemcpy(transferGridDevice, gridHost.baseAddress(), byteSize, cudaMemcpyHostToDevice);
 
         // pad actual grids to avoid edge cell handling
@@ -73 +89 @@
         dim3 gridDim;
         genKernelDimensions(&blockDim, &gridDim);
-        loadGridFromTransferBuffer<<<gridDim, blockDim>>>(curGridDevice, transferGrid, gridWidth() + 2, gridWidth());
+        loadGridFromTransferBuffer<<<gridDim, blockDim>>>(curGridDevice, transferGridDevice, gridWidth() + 2, gridWidth());
+
+        allocInputBuffer(Coord<2>(1, 1));
         checkCudaError();
     }
@@ -79 +97 @@
     virtual ~GPUSimulator()
     {
-        cudaFree(&curGridDevice);
-        cudaFree(&newGridDevice);
+        cudaFree(transferGridDevice);
+        cudaFree(curGridDevice);
+        cudaFree(newGridDevice);
+        cudaFree(inputBufferDevice);
     }
 
@@ -104 +124 @@
         dim3 gridDim;
         genKernelDimensions(&blockDim, &gridDim);
-        storeGridToTransferBuffer<<<gridDim, blockDim>>>(curGridDevice, transferGrid, gridWidth() + 2, gridWidth());
+        storeGridToTransferBuffer<<<gridDim, blockDim>>>(curGridDevice, transferGridDevice, gridWidth() + 2, gridWidth());
         int byteSize = gridHost.getDimensions().prod() * sizeof(CELL_TYPE);
-        cudaMemcpy(gridHost.baseAddress(), transferGrid, byteSize, cudaMemcpyDeviceToHost);
+        cudaMemcpy(gridHost.baseAddress(), transferGridDevice, byteSize, cudaMemcpyDeviceToHost);
         checkCudaError();
         return &gridHost;
+    }
+
+    virtual void updateInputBuffer(unsigned char *data, const Coord<DIM>& inputDim)
+    {
+        allocInputBuffer(inputDim);
+        // fixme: set element size externally
+        int byteSize = inputDim.prod() * 3 * sizeof(unsigned char);
+        cudaMemcpy(inputBufferDevice, data, byteSize, cudaMemcpyHostToDevice);
+  
+        Coord<2> simGridDim = this->getInitializer()->gridDimensions();
+        float factorX = 1.0 * inputDim.x() / simGridDim.x();
+        float factorY = 1.0 * inputDim.y() / simGridDim.y();
+
+        dim3 blockDim;
+        dim3 gridDim;
+        genKernelDimensions(&blockDim, &gridDim);
+        updateCam<<<gridDim, blockDim>>>(curGridDevice, inputBufferDevice, factorX, factorY, gridWidth() + 2, inputDim.x());
     }
 
 private:
     GridType gridHost;
-    CELL_TYPE *transferGrid;
+    CELL_TYPE *transferGridDevice;
     CELL_TYPE *curGridDevice;
     CELL_TYPE *newGridDevice;
-
+    unsigned char *inputBufferDevice;
+    int inputBufferSize;
+
+    void allocInputBuffer(const Coord<DIM>& dim)
+    {
+        // fixme: set element size externally
+        int byteSize = dim.prod() * 3 * sizeof(unsigned char);
+        if (byteSize == inputBufferSize) {
+            return;
+        }
+
+        if (inputBufferDevice != 0) {
+            cudaFree(inputBufferDevice);
+        }
+
+        cudaMalloc(&inputBufferDevice, byteSize);
+        checkCudaError();
+    }
+  
     void checkCudaError()
     {
@@ -159 +214 @@
     typedef Grid<CELL_TYPE, Topology> GridType;
     typedef std::vector<boost::shared_ptr<Writer<CELL_TYPE> > > WriterVector;
+    static const int DIM = Topology::DIMENSIONS;
 
     InteractiveSimulatorGPU(QObject *parent, Initializer<CELL_TYPE> *initializer) :
@@ -170 +226 @@
     virtual void readCam()
     {
-        std::cout << "fixme InteractiveSimulatorGPU::readCam()\n";
+        Coord<DIM> dim = this->initializer->gridDimensions();
+        this->updateInputBuffer(&cameraFrame[0], Coord<2>(cameraFrameWidth, cameraFrameHeight));
     }
 
@@ -182 +239 @@
     virtual void update()
     {
-        // for (int fixme = 0; fixme < 500; ++fixme) 
-        GPUSimulator<CELL_TYPE>::step();
-        // sleep(1);
+        for (int fixme = 0; fixme < 500; ++fixme) 
+            GPUSimulator<CELL_TYPE>::step();
+        sleep(1);
     }
 

src/examples/flowingcanvas/main.cu

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