Merge pull request #857 from UWB-Biocomputing/BenMerge

Moving GPU memory management to OperationManager (clean)

Author: stiber

Simulator/Core/Core.cpp

@@ -206,7 +206,7 @@ int Core::runSimulation(string executableName, string cmdLineArguments)
    }
 
    // Helper function for recorder to register spike history variables for all neurons.
-   simulator.getModel().getLayout().getVertices().registerHistoryVariables();
+   OperationManager::getInstance().executeOperation(Operations::registerHistoryVariables);
 
    // Run simulation
    LOG4CPLUS_TRACE(consoleLogger, "Starting Simulation");
@@ -247,4 +247,4 @@ int Core::runSimulation(string executableName, string cmdLineArguments)
    cout << "time elapsed: " << timeElapsed << endl;
    cout << "ssps (simulation seconds / real time seconds): " << ssps << endl;
    return 0;
-}
+}

Simulator/Core/GPUModel.cpp

@@ -12,6 +12,8 @@
 #include "AllVertices.h"
 #include "Connections.h"
 #include "Global.h"
+#include "OperationManager.h"
+
 #ifdef VALIDATION_MODE
    #include "AllIFNeurons.h"
    #include "OperationManager.h"
@@ -27,51 +29,51 @@ GPUModel::GPUModel() :
    Model::Model(), edgeIndexMapDevice_(nullptr), randNoise_d(nullptr), allVerticesDevice_(nullptr),
    allEdgesDevice_(nullptr)
 {
+   // Register allocNeuronDeviceStruct function as a allocateGPU operation in the OperationManager
+   function<void()> allocateGPU = bind(&GPUModel::allocDeviceStruct, this);
+   OperationManager::getInstance().registerOperation(Operations::allocateGPU, allocateGPU);
+
+   // Register copyCPUtoGPU function as a copyCPUtoGPU operation in the OperationManager
+   function<void()> copyCPUtoGPU = bind(&GPUModel::copyCPUtoGPU, this);
+   OperationManager::getInstance().registerOperation(Operations::copyToGPU, copyCPUtoGPU);
+
+   // Note: We do not register a corresponding copyFromGPU operation here because
+   // we are only copying the synapseIndexMap to the GPU. This map is a read-only lookup table
+   // that gets recreated from scratch on each update. As a result, we only need to allocate,
+   // copy to GPU, and deallocate — there is no meaningful data to copy back from the GPU.
+
+   // Register deleteSynapseImap function as a deallocateGPUMemory operation in the OperationManager
+   function<void()> deallocateGPUMemory = bind(&GPUModel::deleteDeviceStruct, this);
+   OperationManager::getInstance().registerOperation(Operations::deallocateGPUMemory,
+                                                     deallocateGPUMemory);
 }
 
 /// Allocates  and initializes memories on CUDA device.
-/// @param[out] allVerticesDevice          Memory location of the pointer to the vertices list on device memory.
-/// @param[out] allEdgesDevice         Memory location of the pointer to the edges list on device memory.
-void GPUModel::allocDeviceStruct(void **allVerticesDevice, void **allEdgesDevice)
+void GPUModel::allocDeviceStruct()
 {
-   // Get vertices and edges
-   AllVertices &vertices = layout_->getVertices();
-   AllEdges &edges = connections_->getEdges();
-
-   // Allocate vertices and edges structs on GPU device memory
-   vertices.allocVerticesDeviceStruct(allVerticesDevice);
-   edges.allocEdgeDeviceStruct(allEdgesDevice);
-
    // Allocate memory for random noise array
    int numVertices = Simulator::getInstance().getTotalVertices();
    BGSIZE randNoise_d_size = numVertices * sizeof(float);   // size of random noise array
    HANDLE_ERROR(cudaMalloc((void **)&randNoise_d, randNoise_d_size));
 
-   // Copy host vertex and edge arrays into GPU device
-   vertices.copyToDevice(*allVerticesDevice);
-   edges.copyEdgeHostToDevice(*allEdgesDevice);
-
-   // Allocate edge inverse map in device memory
+   // Allocate synapse inverse map in device memory
    allocEdgeIndexMap(numVertices);
 }
 
 /// Copies device memories to host memories and deallocates them.
-/// @param[out] allVerticesDevice          Memory location of the pointer to the vertices list on device memory.
-/// @param[out] allEdgesDevice         Memory location of the pointer to the edges list on device memory.
-void GPUModel::deleteDeviceStruct(void **allVerticesDevice, void **allEdgesDevice)
+void GPUModel::deleteDeviceStruct()
 {
-   // Get vertices and edges
-   AllVertices &vertices = layout_->getVertices();
-   AllEdges &edges = connections_->getEdges();
-
-   // Copy device edge and vertex structs to host memory
-   vertices.copyFromDevice(*allVerticesDevice);
-   // Deallocate device memory
-   vertices.deleteVerticesDeviceStruct(*allVerticesDevice);
-   // Copy device edge and vertex structs to host memory
-   edges.copyEdgeDeviceToHost(*allEdgesDevice);
    // Deallocate device memory
-   edges.deleteEdgeDeviceStruct(*allEdgesDevice);
+   EdgeIndexMapDevice synapseIMapDevice;
+   HANDLE_ERROR(cudaMemcpy(&synapseIMapDevice, edgeIndexMapDevice_, sizeof(EdgeIndexMapDevice),
+                           cudaMemcpyDeviceToHost));
+   HANDLE_ERROR(cudaFree(synapseIMapDevice.outgoingEdgeBegin_));
+   HANDLE_ERROR(cudaFree(synapseIMapDevice.outgoingEdgeCount_));
+   HANDLE_ERROR(cudaFree(synapseIMapDevice.outgoingEdgeIndexMap_));
+   HANDLE_ERROR(cudaFree(synapseIMapDevice.incomingEdgeBegin_));
+   HANDLE_ERROR(cudaFree(synapseIMapDevice.incomingEdgeCount_));
+   HANDLE_ERROR(cudaFree(synapseIMapDevice.incomingEdgeIndexMap_));
+   HANDLE_ERROR(cudaFree(edgeIndexMapDevice_));
    HANDLE_ERROR(cudaFree(randNoise_d));
 }
 
@@ -104,13 +106,9 @@ void GPUModel::setupSim()
    t_gpu_advanceSynapses = 0.0;
    t_gpu_calcSummation = 0.0;
 #endif   // PERFORMANCE_METRICS
-
-   // allocates memories on CUDA device
-   allocDeviceStruct((void **)&allVerticesDevice_, (void **)&allEdgesDevice_);
-
-   EdgeIndexMap &edgeIndexMap = connections_->getEdgeIndexMap();
-   // copy inverse map to the device memory
-   copyEdgeIndexMapHostToDevice(edgeIndexMap, Simulator::getInstance().getTotalVertices());
+   // Allocate and copy neuron/synapse data structures to GPU memory
+   OperationManager::getInstance().executeOperation(Operations::allocateGPU);
+   OperationManager::getInstance().executeOperation(Operations::copyToGPU);
 
    AllEdges &edges = connections_->getEdges();
    // set some parameters used for advanceVerticesDevice
@@ -123,9 +121,10 @@ void GPUModel::setupSim()
 /// Performs any finalization tasks on network following a simulation.
 void GPUModel::finish()
 {
+   // copy device synapse and neuron structs to host memory
+   OperationManager::getInstance().executeOperation(Operations::copyFromGPU);
    // deallocates memories on CUDA device
-   deleteDeviceStruct((void **)&allVerticesDevice_, (void **)&allEdgesDevice_);
-   deleteEdgeIndexMap();
+   OperationManager::getInstance().executeOperation(Operations::deallocateGPUMemory);
 
 #ifdef PERFORMANCE_METRICS
    cudaEventDestroy(start);
@@ -234,7 +233,7 @@ void GPUModel::updateConnections()
    AllVertices &vertices = layout_->getVertices();
    AllEdges &edges = connections_->getEdges();
 
-   vertices.copyFromDevice(allVerticesDevice_);
+   vertices.copyFromDevice();
 
    // Update Connections data
    if (connections_->updateConnections(vertices)) {
@@ -243,8 +242,7 @@ void GPUModel::updateConnections()
       // create edge index map
       connections_->createEdgeIndexMap();
       // copy index map to the device memory
-      copyEdgeIndexMapHostToDevice(connections_->getEdgeIndexMap(),
-                                   Simulator::getInstance().getTotalVertices());
+      copyCPUtoGPU();
    }
 }
 
@@ -281,83 +279,67 @@ void GPUModel::allocEdgeIndexMap(int count)
                            cudaMemcpyHostToDevice));
 }
 
-/// Deallocate device memory for edge inverse map.
-void GPUModel::deleteEdgeIndexMap()
-{
-   EdgeIndexMapDevice edgeIndexMapDevice;
-   HANDLE_ERROR(cudaMemcpy(&edgeIndexMapDevice, edgeIndexMapDevice_, sizeof(EdgeIndexMapDevice),
-                           cudaMemcpyDeviceToHost));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice.outgoingEdgeBegin_));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice.outgoingEdgeCount_));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice.outgoingEdgeIndexMap_));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice.incomingEdgeBegin_));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice.incomingEdgeCount_));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice.incomingEdgeIndexMap_));
-   HANDLE_ERROR(cudaFree(edgeIndexMapDevice_));
-}
-
-/// Copy EdgeIndexMap in host memory to EdgeIndexMap in device memory.
-/// @param  edgeIndexMapHost		Reference to the EdgeIndexMap in host memory.
-void GPUModel::copyEdgeIndexMapHostToDevice(EdgeIndexMap &edgeIndexMapHost, int numVertices)
+/// Allocate and Copy CPU Synapse data to GPU.
+void GPUModel::copyCPUtoGPU()
 {
-   AllEdges &edges = connections_->getEdges();
-   int totalEdgeCount = edges.totalEdgeCount_;
-   if (totalEdgeCount == 0)
+   EdgeIndexMap synapseIndexMapHost = connections_->getEdgeIndexMap();
+   int numVertices = Simulator::getInstance().getTotalVertices();
+   AllEdges &synapses = connections_->getEdges();
+   int totalSynapseCount = dynamic_cast<AllEdges &>(synapses).totalEdgeCount_;
+   if (totalSynapseCount == 0)
       return;
-   EdgeIndexMapDevice edgeIndexMapDevice;
-   HANDLE_ERROR(cudaMemcpy(&edgeIndexMapDevice, edgeIndexMapDevice_, sizeof(EdgeIndexMapDevice),
+   EdgeIndexMapDevice synapseIMapDevice;
+   HANDLE_ERROR(cudaMemcpy(&synapseIMapDevice, edgeIndexMapDevice_, sizeof(EdgeIndexMapDevice),
                            cudaMemcpyDeviceToHost));
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice.outgoingEdgeBegin_,
-                           edgeIndexMapHost.outgoingEdgeBegin_.data(), numVertices * sizeof(BGSIZE),
-                           cudaMemcpyHostToDevice));
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice.outgoingEdgeCount_,
-                           edgeIndexMapHost.outgoingEdgeCount_.data(), numVertices * sizeof(BGSIZE),
-                           cudaMemcpyHostToDevice));
-   if (edgeIndexMapDevice.outgoingEdgeIndexMap_ != nullptr) {
-      HANDLE_ERROR(cudaFree(edgeIndexMapDevice.outgoingEdgeIndexMap_));
+   HANDLE_ERROR(cudaMemcpy(synapseIMapDevice.outgoingEdgeBegin_,
+                           synapseIndexMapHost.outgoingEdgeBegin_.data(),
+                           numVertices * sizeof(BGSIZE), cudaMemcpyHostToDevice));
+   HANDLE_ERROR(cudaMemcpy(synapseIMapDevice.outgoingEdgeCount_,
+                           synapseIndexMapHost.outgoingEdgeCount_.data(),
+                           numVertices * sizeof(BGSIZE), cudaMemcpyHostToDevice));
+   if (synapseIMapDevice.outgoingEdgeIndexMap_ != nullptr) {
+      HANDLE_ERROR(cudaFree(synapseIMapDevice.outgoingEdgeIndexMap_));
    }
-   HANDLE_ERROR(cudaMalloc((void **)&edgeIndexMapDevice.outgoingEdgeIndexMap_,
-                           totalEdgeCount * sizeof(BGSIZE)));
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice.outgoingEdgeIndexMap_,
-                           edgeIndexMapHost.outgoingEdgeIndexMap_.data(),
-                           totalEdgeCount * sizeof(BGSIZE), cudaMemcpyHostToDevice));
+   HANDLE_ERROR(cudaMalloc((void **)&synapseIMapDevice.outgoingEdgeIndexMap_,
+                           totalSynapseCount * sizeof(BGSIZE)));
+   HANDLE_ERROR(cudaMemcpy(synapseIMapDevice.outgoingEdgeIndexMap_,
+                           synapseIndexMapHost.outgoingEdgeIndexMap_.data(),
+                           totalSynapseCount * sizeof(BGSIZE), cudaMemcpyHostToDevice));
    // active synapse map
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice.incomingEdgeBegin_,
-                           edgeIndexMapHost.incomingEdgeBegin_.data(), numVertices * sizeof(BGSIZE),
-                           cudaMemcpyHostToDevice));
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice.incomingEdgeCount_,
-                           edgeIndexMapHost.incomingEdgeCount_.data(), numVertices * sizeof(BGSIZE),
-                           cudaMemcpyHostToDevice));
+   HANDLE_ERROR(cudaMemcpy(synapseIMapDevice.incomingEdgeBegin_,
+                           synapseIndexMapHost.incomingEdgeBegin_.data(),
+                           numVertices * sizeof(BGSIZE), cudaMemcpyHostToDevice));
+   HANDLE_ERROR(cudaMemcpy(synapseIMapDevice.incomingEdgeCount_,
+                           synapseIndexMapHost.incomingEdgeCount_.data(),
+                           numVertices * sizeof(BGSIZE), cudaMemcpyHostToDevice));
    // the number of synapses may change, so we reallocate the memory
-   if (edgeIndexMapDevice.incomingEdgeIndexMap_ != nullptr) {
-      HANDLE_ERROR(cudaFree(edgeIndexMapDevice.incomingEdgeIndexMap_));
-      edgeIndexMapDevice.incomingEdgeIndexMap_ = nullptr;
+   if (synapseIMapDevice.incomingEdgeIndexMap_ != nullptr) {
+      HANDLE_ERROR(cudaFree(synapseIMapDevice.incomingEdgeIndexMap_));
+      synapseIMapDevice.incomingEdgeIndexMap_ = nullptr;
    }
-   HANDLE_ERROR(cudaMalloc((void **)&edgeIndexMapDevice.incomingEdgeIndexMap_,
-                           totalEdgeCount * sizeof(BGSIZE)));
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice.incomingEdgeIndexMap_,
-                           edgeIndexMapHost.incomingEdgeIndexMap_.data(),
-                           totalEdgeCount * sizeof(BGSIZE), cudaMemcpyHostToDevice));
-   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice_, &edgeIndexMapDevice, sizeof(EdgeIndexMapDevice),
+   HANDLE_ERROR(cudaMalloc((void **)&synapseIMapDevice.incomingEdgeIndexMap_,
+                           totalSynapseCount * sizeof(BGSIZE)));
+   HANDLE_ERROR(cudaMemcpy(synapseIMapDevice.incomingEdgeIndexMap_,
+                           synapseIndexMapHost.incomingEdgeIndexMap_.data(),
+                           totalSynapseCount * sizeof(BGSIZE), cudaMemcpyHostToDevice));
+   HANDLE_ERROR(cudaMemcpy(edgeIndexMapDevice_, &synapseIMapDevice, sizeof(EdgeIndexMapDevice),
                            cudaMemcpyHostToDevice));
 }
 
-/// Copy GPU edge data to CPU.
-void GPUModel::copyGPUtoCPU()
+/// Print out EdgeProps on the GPU.
+void GPUModel::printGPUEdgesPropsModel() const
 {
-   // copy device edge structs to host memory
-   connections_->getEdges().copyEdgeDeviceToHost(allEdgesDevice_);
+   connections_->getEdges().printGPUEdgesProps(allEdgesDevice_);
 }
 
-/// Copy CPU edge data to GPU.
-void GPUModel::copyCPUtoGPU()
+/// Getter for neuron structure in device memory
+AllVerticesDeviceProperties *&GPUModel::getAllVerticesDevice()
 {
-   // copy host edge structs to device memory
-   connections_->getEdges().copyEdgeHostToDevice(allEdgesDevice_);
+   return allVerticesDevice_;
 }
 
-/// Print out EdgeProps on the GPU.
-void GPUModel::printGPUEdgesPropsModel() const
+/// Getter for synapse structures in device memory
+AllEdgesDeviceProperties *&GPUModel::getAllEdgesDevice()
 {
-   connections_->getEdges().printGPUEdgesProps(allEdgesDevice_);
-}
+   return allEdgesDevice_;
+}

Simulator/Core/GPUModel.h

@@ -22,7 +22,10 @@
 #pragma once
 
 #include "AllEdges.h"
+#include "AllSpikingNeurons.h"
+#include "AllSpikingSynapses.h"
 #include "AllVertices.h"
+#include "OperationManager.h"
 
 #ifdef VALIDATION_MODE
    #include <fstream>
@@ -81,25 +84,33 @@ class GPUModel : public Model {
    /// over the past epoch. Should be called once every epoch.
    virtual void updateConnections() override;
 
-   /// Copy GPU edge data to CPU.
-   virtual void copyGPUtoCPU() override;
-
-   /// Copy CPU edge data to GPU.
+   /// Copies neuron and synapse data from CPU to GPU memory.
+   /// TODO: Refactor this. Currently, GPUModel handles low-level memory transfer for vertices and edges.
+   ///       Consider moving this responsibility to a more appropriate class, such as a dedicated memory manager
+   ///       or the OperationManager, to better separate concerns and keep the model focused on high-level coordination.
    virtual void copyCPUtoGPU() override;
 
+   // GPUModel itself does not have anything to be copied back, this function is a
+   // dummy function just to make GPUModel non virtual
+   virtual void copyGPUtoCPU() override
+   {
+   }
+
    /// Print out EdgeProps on the GPU.
    void printGPUEdgesPropsModel() const;
 
+   /// Getter for edge (synapse) structures in device memory
+   AllEdgesDeviceProperties *&getAllEdgesDevice();
+
+   /// Getter for vertex (neuron) structures in device memory
+   AllVerticesDeviceProperties *&getAllVerticesDevice();
+
 protected:
    /// Allocates  and initializes memories on CUDA device.
-   /// @param[out] allVerticesDevice          Memory location of the pointer to the vertices list on device memory.
-   /// @param[out] allEdgesDevice         Memory location of the pointer to the edges list on device memory.
-   void allocDeviceStruct(void **allVerticesDevice, void **allEdgesDevice);
+   void allocDeviceStruct();
 
-   /// Copies device memories to host memories and deallocates them.
-   /// @param[out] allVerticesDevice          Memory location of the pointer to the vertices list on device memory.
-   /// @param[out] allEdgesDevice         Memory location of the pointer to the edges list on device memory.
-   virtual void deleteDeviceStruct(void **allVerticesDevice, void **allEdgesDevice);
+   /// Deallocates device memories.
+   virtual void deleteDeviceStruct();
 
    /// Pointer to device random noise array.
    float *randNoise_d;
@@ -118,11 +129,6 @@ class GPUModel : public Model {
 private:
    void allocEdgeIndexMap(int count);
 
-   void deleteEdgeIndexMap();
-
-public:   //2020/03/14 changed to public for accessing in Core
-   void copyEdgeIndexMapHostToDevice(EdgeIndexMap &edgeIndexMapHost, int numVertices);
-
 private:
    void updateHistory();
 
@@ -144,4 +150,4 @@ void normalMTGPU(float *randNoise_d);
 void initMTGPU(unsigned int seed, unsigned int blocks, unsigned int threads, unsigned int nPerRng,
                unsigned int mt_rng_count);
 }
-#endif
+#endif

Simulator/Core/OperationManager.cpp

@@ -71,7 +71,9 @@ string OperationManager::operationToString(const Operations &operation) const
          return "copyToGPU";
       case Operations::copyFromGPU:
          return "copyFromGPU";
+      case Operations::allocateGPU:
+         return "allocateGPU";
       default:
          return "Operation isn't in OperationManager::operationToString()";
    }
-}
+}

Simulator/Core/Operations.h

@@ -20,5 +20,7 @@ enum class Operations {
    deallocateGPUMemory,   // Make sure deallocate memory isn't called until all GPU memory is copied back.
    restoreToDefault,   // Not sure what this refers to.
    copyToGPU,
-   copyFromGPU
+   copyFromGPU,
+   allocateGPU,
+   registerHistoryVariables
 };