[issue-858] GPU implementation of 911 vertices #880

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft

NicolasJPosey wants to merge 67 commits into PoseyDevelopment from issue-858-911vertices-gpu-implementation

CMakeLists.txt

-Original file line number
+Diff line change
@@ Expand Up / @@ -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()
@@ Expand Down Expand Up / @@ -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})
@@ Expand Down @@

Simulator/Connections/NG911/Connections911.cpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -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)
     {
@@ Expand Down @@

Simulator/Connections/NG911/Connections911.h

-Original file line number
+Diff line change
@@ Expand Up / @@ -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);
@@ Expand Down @@

Simulator/Core/FunctionNodes/GenericFunctionNode.h

-Original file line number
+Diff line change
@@ Expand Up / @@ -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.
@@ Expand Down @@

Simulator/Core/FunctionNodes/IFunctionNode.h

-Original file line number
+Diff line change
@@ Expand Up / @@ -10,6 +10,7 @@ @@
     #pragma once
     #include "Operations.h"
+    #include <cstdint>   ///for uint64_t
     using namespace std;
@@ Expand All / @@ -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_;
@@ Expand Down @@

Simulator/Core/FunctionNodes/TwoUint64ArgFunctionNode.cpp

-Original file line number
+Diff line change
@@ -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

-Original file line number
+Diff line change
@@ -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

-Original file line number
+Diff line change
@@ Expand Up / @@ -12,6 +12,7 @@ @@
     #include "AllVertices.h"
     #include "Connections.h"
     #include "Global.h"
+    #include "MersenneTwister_d.h"
     #include "OperationManager.h"
     #ifdef VALIDATION_MODE
@@ Expand Down Expand Up / @@ -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.
@@ Expand Down Expand Up / @@ -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);
@@ Expand Down Expand Up / @@ -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,
@@ Expand Down Expand Up / @@ -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_);
@@ Expand Down Expand Up @@
     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

-Original file line number
+Diff line change
@@ Expand Up / @@ -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;
@@ Expand Down @@

Simulator/Core/OperationManager.cpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -13,6 +13,7 @@ @@
     #include "OperationManager.h"
     #include "GenericFunctionNode.h"
+    #include "TwoUint64ArgFunctionNode.h"
     #include <list>
     #include <memory>
     #include <string>
@@ Expand Down Expand Up @@
        }
     }
+    /// @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
           }
        }
     }
@@ Expand All @@
              return "copyFromGPU";
           case Operations::allocateGPU:
              return "allocateGPU";
+          case Operations::loadEpochInputs:
+             return "loadEpochInputs";
           default:
              return "Operation isn't in OperationManager::operationToString()";
        }
@@ Expand Down @@

Simulator/Core/OperationManager.h

-Original file line number
+Diff line change
@@ Expand Up / @@ -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;
@@ Expand Down @@

Simulator/Core/Operations.h

-Original file line number
+Diff line change
@@ Expand Up / @@ -22,5 +22,6 @@ enum class Operations { @@
        copyToGPU,
        copyFromGPU,
        allocateGPU,
-       registerHistoryVariables
+       registerHistoryVariables,
+       loadEpochInputs
     };

Simulator/Core/Simulator.cpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -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) {
@@ Expand Down @@

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[issue-858] GPU implementation of 911 vertices #880

Uh oh!

Diff view

Diff view

There are no files selected for viewing

Uh oh!

[issue-858] GPU implementation of 911 vertices #880

Are you sure you want to change the base?

Uh oh!

[issue-858] GPU implementation of 911 vertices #880

Uh oh!

Uh oh!

Diff view

Diff view

There are no files selected for viewing

Uh oh!