[issue-858] GPU implementation of 911 vertices

CMakeLists.txt

@@ -264,7 +264,8 @@ if(ENABLE_CUDA)
         Simulator/Vertices/Neuro/AllLIFNeurons_d.cpp
         Simulator/Vertices/Neuro/AllIZHNeurons_d.cpp
         Simulator/Vertices/Neuro/AllIFNeurons_d.cpp
-        Simulator/Vertices/Neuro/AllSpikingNeurons_d.cpp)
+        Simulator/Vertices/Neuro/AllSpikingNeurons_d.cpp
+        Simulator/Vertices/NG911/All911Vertices_d.cpp)
         set_source_files_properties(${cuda_VerticesSources} PROPERTIES LANGUAGE CUDA)
 endif()
 
@@ -296,6 +297,7 @@ else()
         list(REMOVE_ITEM Vertices_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Vertices/Neuro/AllIZHNeurons_d.cpp")
         list(REMOVE_ITEM Vertices_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Vertices/Neuro/AllIFNeurons_d.cpp")
         list(REMOVE_ITEM Vertices_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Vertices/Neuro/AllSpikingNeurons_d.cpp")
+        list(REMOVE_ITEM Vertices_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Vertices/NG911/All911Vertices_d.cpp")
 
         add_library(Vertices STATIC ${Vertices_Source})
 

Simulator/Connections/NG911/Connections911.cpp

@@ -96,52 +96,6 @@ bool Connections911::updateConnections()
 }
 
 
-/// Finds the outgoing edge from the given vertex to the Responder closest to
-/// the emergency call location
-BGSIZE Connections911::getEdgeToClosestResponder(const Call &call, BGSIZE vertexIdx)
-{
-   All911Edges &edges911 = dynamic_cast<All911Edges &>(*edges_);
-
-   vertexType requiredType;
-   if (call.type == "Law")
-      requiredType = vertexType::LAW;
-   else if (call.type == "EMS")
-      requiredType = vertexType::EMS;
-   else if (call.type == "Fire")
-      requiredType = vertexType::FIRE;
-
-   // loop over the outgoing edges looking for the responder with the shortest
-   // Euclidean distance to the call's location.
-   BGSIZE startOutEdg = synapseIndexMap_->outgoingEdgeBegin_[vertexIdx];
-   BGSIZE outEdgCount = synapseIndexMap_->outgoingEdgeCount_[vertexIdx];
-   Layout911 &layout911
-      = dynamic_cast<Layout911 &>(Simulator::getInstance().getModel().getLayout());
-
-   BGSIZE resp, respEdge;
-   double minDistance = numeric_limits<double>::max();
-   for (BGSIZE eIdxMap = startOutEdg; eIdxMap < startOutEdg + outEdgCount; ++eIdxMap) {
-      BGSIZE outEdg = synapseIndexMap_->outgoingEdgeIndexMap_[eIdxMap];
-      assert(edges911.inUse_[outEdg]);   // Edge must be in use
-
-      BGSIZE dstVertex = edges911.destVertexIndex_[outEdg];
-      if (layout911.vertexTypeMap_[dstVertex] == requiredType) {
-         double distance = layout911.getDistance(dstVertex, call.x, call.y);
-
-         if (distance < minDistance) {
-            minDistance = distance;
-            resp = dstVertex;
-            respEdge = outEdg;
-         }
-      }
-   }
-
-   // We must have found the closest responder of the right type
-   assert(minDistance < numeric_limits<double>::max());
-   assert(layout911.vertexTypeMap_[resp] == requiredType);
-   return respEdge;
-}
-
-
 ///  Randomly delete 1 PSAP and rewire all the edges around it.
 bool Connections911::erasePSAP(AllVertices &vertices, Layout &layout)
 {

Simulator/Connections/NG911/Connections911.h

@@ -79,14 +79,6 @@ class Connections911 : public Connections {
    ///  @return true if successful, false otherwise.
    virtual bool updateConnections() override;
 
-   /// Finds the outgoing edge from the given vertex to the Responder closest to
-   /// the emergency call location
-   ///
-   /// @param call         The call that needs a Responder
-   /// @param vertexIdx    The index of the vertex serving the call (A PSAP)
-   /// @return    The index of the outgoing edge to the closest Responder
-   BGSIZE getEdgeToClosestResponder(const Call &call, BGSIZE vertexIdx);
-
    ///  Returns the complete list of all deleted or added edges as a string.
    ///  @return xml representation of all deleted or added edges
    string changedEdgesToXML(bool added);

Simulator/Core/FunctionNodes/GenericFunctionNode.h

@@ -23,7 +23,13 @@ class GenericFunctionNode : public IFunctionNode {
    ~GenericFunctionNode() = default;
 
    /// Invokes the stored function if the sent operation type matches the operation type the function is stored as.
-   bool invokeFunction(const Operations &operation) const override;
+   virtual bool invokeFunction(const Operations &operation) const override;
+
+   /// TODO: Remove when IFunctionNode supports functions with non-empty signatures
+   virtual bool invokeFunction(const Operations &operation, uint64_t arg1, uint64_t arg2) const
+   {
+      return false;
+   }
 
 private:
    std::function<void()> function_;   ///< Stored function.

Simulator/Core/FunctionNodes/IFunctionNode.h

@@ -10,6 +10,7 @@
 #pragma once
 
 #include "Operations.h"
+#include <cstdint>   ///for uint64_t
 
 using namespace std;
 
@@ -18,9 +19,14 @@ class IFunctionNode {
    /// Destructor.
    virtual ~IFunctionNode() = default;
 
+   /// TODO: Need to refactor to allow for passing in arguments. Otherwise, FunctionNode classes can not support
+   /// non-empty signatures.
    /// Invokes the stored function if the sent operation type matches the operation type the function is stored as.
    virtual bool invokeFunction(const Operations &operation) const = 0;
 
+   /// Invokes the stored function using the two arguments as input
+   virtual bool invokeFunction(const Operations &operation, uint64_t arg1, uint64_t arg2) const = 0;
+
 protected:
    /// The operation type of the stored function.
    Operations operationType_;

Simulator/Core/FunctionNodes/TwoUint64ArgFunctionNode.cpp

@@ -0,0 +1,33 @@
+/**
+ *  @file TwoUint64ArgFunctionNode.cpp
+ * 
+ *  @ingroup Simulator/Core/FunctionNodes
+ *
+ *  @brief Stores a function with two uint64_t args to invoke. Used by operation manager to store functions to defined by an operation type.
+ * 
+ *  Function Signature supported : void (uint64_t,uint64_t)
+ *
+ */
+
+#include "TwoUint64ArgFunctionNode.h"
+#include "Operations.h"
+#include <functional>
+
+/// Constructor, Function Signature: void (uint64_t, uint64_t)
+TwoUint64ArgFunctionNode::TwoUint64ArgFunctionNode(
+   const Operations &operation, const std::function<void(uint64_t, uint64_t)> &func)
+{
+   operationType_ = operation;
+   function_ = func;
+}
+
+/// Invokes the stored function if the sent operation type matches the operation type the function is stored as.
+bool TwoUint64ArgFunctionNode::invokeFunction(const Operations &operation, uint64_t arg1,
+                                              uint64_t arg2) const
+{
+   if (operation == operationType_) {
+      __invoke(function_, arg1, arg2);
+      return true;
+   }
+   return false;
+}

Simulator/Core/FunctionNodes/TwoUint64ArgFunctionNode.h

@@ -0,0 +1,38 @@
+/**
+ *  @file TwoUint64ArgFunctionNode.h
+ * 
+ *  @ingroup Simulator/Core/FunctionNodes
+ *
+ *  @brief Stores a function with two uint64_t args to invoke. Used by operation manager to store functions to defined by an operation type.
+ *
+ */
+
+#pragma once
+
+#include "IFunctionNode.h"
+#include <functional>
+
+using namespace std;
+
+class TwoUint64ArgFunctionNode : public IFunctionNode {
+public:
+   /// Constructor, Function Signature: void ()
+   TwoUint64ArgFunctionNode(const Operations &operationType,
+                            const std::function<void(uint64_t, uint64_t)> &function);
+
+   /// Destructor
+   ~TwoUint64ArgFunctionNode() = default;
+
+   /// TODO: Remove when IFunctionNode supports functions with non-empty signatures
+   virtual bool invokeFunction(const Operations &operation) const
+   {
+      return false;
+   }
+
+   /// Invokes the stored function if the sent operation type matches the operation type the function is stored as.
+   virtual bool invokeFunction(const Operations &operation, uint64_t arg1,
+                               uint64_t arg2) const override;
+
+private:
+   std::function<void(uint64_t, uint64_t)> function_;   ///< Stored function.
+};

Simulator/Core/GPUModel.cpp

@@ -12,6 +12,7 @@
 #include "AllVertices.h"
 #include "Connections.h"
 #include "Global.h"
+#include "MersenneTwister_d.h"
 #include "OperationManager.h"
 
 #ifdef VALIDATION_MODE
@@ -52,12 +53,15 @@ GPUModel::GPUModel() :
 void GPUModel::allocDeviceStruct()
 {
    // Allocate memory for random noise array
-   int numVertices = Simulator::getInstance().getTotalVertices();
-   BGSIZE randNoise_d_size = numVertices * sizeof(float);   // size of random noise array
+   int numVerticesNeedingNoise = layout_->getVertices().getNumberOfVerticesNeedingDeviceNoise();
+   int numberOfNoiseElements = roundUpNumberOfNoiseElements(numVerticesNeedingNoise);
+   LOG4CPLUS_DEBUG(fileLogger_,
+                   "Number of elements allocated for noise: " << numberOfNoiseElements);
+   BGSIZE randNoise_d_size = numberOfNoiseElements * sizeof(float);   // size of random noise array
    HANDLE_ERROR(cudaMalloc((void **)&randNoise_d, randNoise_d_size));
 
    // Allocate synapse inverse map in device memory
-   allocEdgeIndexMap(numVertices);
+   allocEdgeIndexMap(Simulator::getInstance().getTotalVertices());
 }
 
 /// Copies device memories to host memories and deallocates them.
@@ -91,9 +95,20 @@ void GPUModel::setupSim()
    int rng_blocks = 25;   //# of blocks the kernel will use
    int rng_nPerRng
       = 4;   //# of iterations per thread (thread granularity, # of rands generated per thread)
-   int rng_mt_rng_count = Simulator::getInstance().getTotalVertices()
-                          / rng_nPerRng;   //# of threads to generate for numVertices rand #s
+   int numVerticesNeedingNoise = layout_->getVertices().getNumberOfVerticesNeedingDeviceNoise();
+   int numberOfNoiseElements = roundUpNumberOfNoiseElements(numVerticesNeedingNoise);
+   int rng_mt_rng_count
+      = numberOfNoiseElements / rng_nPerRng;   //# of threads to generate for numVertices rand #s
+   assert(rng_mt_rng_count <= MT_RNG_COUNT);
    int rng_threads = rng_mt_rng_count / rng_blocks;   //# threads per block needed
+   LOG4CPLUS_DEBUG(fileLogger_, "initMTGPU state: " << endl
+                                                    << "Noise seed: "
+                                                    << Simulator::getInstance().getNoiseRngSeed()
+                                                    << endl
+                                                    << "RNG_blocks: " << rng_blocks << endl
+                                                    << "RNG_threads: " << rng_threads << endl
+                                                    << "RNG_nPerRng: " << rng_nPerRng << endl
+                                                    << "Count: " << rng_mt_rng_count);
    initMTGPU(Simulator::getInstance().getNoiseRngSeed(), rng_blocks, rng_threads, rng_nPerRng,
              rng_mt_rng_count);
 
@@ -165,7 +180,6 @@ void GPUModel::advance()
    cudaLapTime(t_gpu_rndGeneration);
    cudaStartTimer();
 #endif   // PERFORMANCE_METRICS
-
    // display running info to console
    // Advance vertices ------------->
    vertices.advanceVertices(edges, allVerticesDevice_, allEdgesDevice_, randNoise_d,
@@ -217,7 +231,6 @@ void GPUModel::advance()
    cudaLapTime(t_gpu_advanceSynapses);
    cudaStartTimer();
 #endif   // PERFORMANCE_METRICS
-
    // integrate the inputs of the vertices
    vertices.integrateVertexInputs(allVerticesDevice_, edgeIndexMapDevice_, allEdgesDevice_);
 
@@ -342,4 +355,16 @@ AllVerticesDeviceProperties *&GPUModel::getAllVerticesDevice()
 AllEdgesDeviceProperties *&GPUModel::getAllEdgesDevice()
 {
    return allEdgesDevice_;
+}
+
+int GPUModel::roundUpNumberOfNoiseElements(int input)
+{
+   // MersenneTwister requires the number of elements to be 100 or more and a multiple of 100
+   // To deal with this, we take the input and round it up to the nearest multiple of 100.
+   assert(input > 0);
+   // Already a multiple of 100 so return
+   if (input % 100 == 0)
+      return input;
+   // Return the next highest multiple of 100
+   return ((input + 99) / 100) * 100;
 }

Simulator/Core/GPUModel.h

@@ -112,6 +112,10 @@ class GPUModel : public Model {
    /// Deallocates device memories.
    virtual void deleteDeviceStruct();
 
+   /// Takes the input and returns a rounded up number of elements to
+   /// use for generating device noise.
+   int roundUpNumberOfNoiseElements(int input);
+
    /// Pointer to device random noise array.
    float *randNoise_d;
 

Simulator/Core/OperationManager.cpp

@@ -13,6 +13,7 @@
 
 #include "OperationManager.h"
 #include "GenericFunctionNode.h"
+#include "TwoUint64ArgFunctionNode.h"
 #include <list>
 #include <memory>
 #include <string>
@@ -40,13 +41,44 @@ void OperationManager::registerOperation(const Operations &operation,
    }
 }
 
+/// @brief Handles function signature: void (uint64_t,uint64_t).
+/// @param operation The Operation type that will use the input function.
+/// @param function The function invoked for the operation. Takes in two arguments of type uint64_t
+void OperationManager::registerOperation(const Operations &operation,
+                                         const function<void(uint64_t, uint64_t)> &function)
+{
+   try {
+      functionList_.push_back(
+         unique_ptr<IFunctionNode>(new TwoUint64ArgFunctionNode(operation, function)));
+   } catch (exception e) {
+      LOG4CPLUS_FATAL(logger_, string(e.what())
+                                  + ". Push back failed in OperationManager::registerOperation");
+      throw runtime_error(string(e.what()) + " in OperationManager::registerOperation");
+   }
+}
+
 /// Takes in a operation type and invokes all registered functions that are classified as that operation type.
 void OperationManager::executeOperation(const Operations &operation) const
 {
    LOG4CPLUS_INFO(logger_, "Executing operation " + operationToString(operation));
    if (functionList_.size() > 0) {
       for (auto i = functionList_.begin(); i != functionList_.end(); ++i) {
          (*i)->invokeFunction(operation);
+         //TODO: Throw fatal if false
+      }
+   }
+}
+
+/// Take in a operation type and invokes all registered functions that are classified as that operation type using the input arguments.
+void OperationManager::executeOperation(const Operations &operation, uint64_t arg1,
+                                        uint64_t arg2) const
+{
+   LOG4CPLUS_INFO(logger_, "Executing operation " + operationToString(operation));
+   /// TODO: Should we check anything about arg1 and arg2 before passing to the invoke???
+   if (functionList_.size() > 0) {
+      for (auto i = functionList_.begin(); i != functionList_.end(); ++i) {
+         (*i)->invokeFunction(operation, arg1, arg2);
+         //TODO: Throw fatal if false
       }
    }
 }
@@ -73,6 +105,8 @@ string OperationManager::operationToString(const Operations &operation) const
          return "copyFromGPU";
       case Operations::allocateGPU:
          return "allocateGPU";
+      case Operations::loadEpochInputs:
+         return "loadEpochInputs";
       default:
          return "Operation isn't in OperationManager::operationToString()";
    }

Simulator/Core/OperationManager.h

@@ -35,9 +35,16 @@ class OperationManager {
    /// Handles function signature: void ()
    void registerOperation(const Operations &operation, const function<void()> &function);
 
+   /// Handles function signature: void (uint64_t,uint64_t)
+   void registerOperation(const Operations &operation,
+                          const function<void(uint64_t, uint64_t)> &function);
+
    /// Takes in a operation type and invokes all registered functions that are classified as that operation type.
    void executeOperation(const Operations &operation) const;
 
+   /// Take in a operation type and invokes all registered functions that are classified as that operation type using the input arguments.
+   void executeOperation(const Operations &operation, uint64_t arg1, uint64_t arg2) const;
+
    /// Takes in the operation enum and returns the enum as a string. Used for debugging purposes.
    string operationToString(const Operations &operation) const;
 

Simulator/Core/Operations.h

@@ -22,5 +22,6 @@ enum class Operations {
    copyToGPU,
    copyFromGPU,
    allocateGPU,
-   registerHistoryVariables
+   registerHistoryVariables,
+   loadEpochInputs
 };

Simulator/Core/Simulator.cpp

@@ -188,7 +188,8 @@ void Simulator::advanceEpoch(int currentEpoch) const
    uint64_t count = 0;
    // Compute step number at end of this simulation epoch
    uint64_t endStep = g_simulationStep + static_cast<uint64_t>(epochDuration_ / deltaT_);
-   model_->getLayout().getVertices().loadEpochInputs(g_simulationStep, endStep);
+   OperationManager::getInstance().executeOperation(Operations::loadEpochInputs, g_simulationStep,
+                                                    endStep);
    // DEBUG_MID(model->logSimStep();) // Generic model debug call
    uint64_t onePercent = (epochDuration_ / deltaT_) * numEpochs_ * 0.01;
    while (g_simulationStep < endStep) {