Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 6 additions & 2 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -255,15 +255,17 @@ if(ENABLE_CUDA)
Simulator/Edges/Neuro/AllSpikingSynapses_d.cpp
Simulator/Edges/Neuro/AllSTDPSynapses_d.cpp
Simulator/Edges/Neuro/AllDSSynapses_d.cpp
Simulator/Edges/Neuro/AllDynamicSTDPSynapses_d.cpp)
Simulator/Edges/Neuro/AllDynamicSTDPSynapses_d.cpp
Simulator/Edges/NG911/All911Edges_d.cpp)
set_source_files_properties(${cuda_EdgesSources} PROPERTIES LANGUAGE CUDA)

set(cuda_VerticesSources
Simulator/Vertices/Neuro/AllVerticesDeviceFuncs_d.cpp
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 @@ -295,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 All @@ -313,6 +316,7 @@ else()
list(REMOVE_ITEM Edges_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Edges/Neuro/AllSTDPSynapses_d.cpp")
list(REMOVE_ITEM Edges_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Edges/Neuro/AllDSSynapses_d.cpp")
list(REMOVE_ITEM Edges_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Edges/Neuro/AllDynamicSTDPSynapses_d.cpp")
list(REMOVE_ITEM Edges_Source "${CMAKE_CURRENT_SOURCE_DIR}/Simulator/Edges/NG911/All911Edges_d.cpp")

add_library(Edges STATIC ${Edges_Source})

Expand Down
46 changes: 0 additions & 46 deletions Simulator/Connections/NG911/Connections911.cpp
Original file line number Diff 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
8 changes: 0 additions & 8 deletions Simulator/Connections/NG911/Connections911.h
Original file line number Diff 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
8 changes: 7 additions & 1 deletion Simulator/Core/FunctionNodes/GenericFunctionNode.h
Original file line number Diff 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
6 changes: 6 additions & 0 deletions Simulator/Core/FunctionNodes/IFunctionNode.h
Original file line number Diff 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
33 changes: 33 additions & 0 deletions Simulator/Core/FunctionNodes/TwoUint64ArgFunctionNode.cpp
Original file line number Diff 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;
}
38 changes: 38 additions & 0 deletions Simulator/Core/FunctionNodes/TwoUint64ArgFunctionNode.h
Original file line number Diff 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.
};
39 changes: 32 additions & 7 deletions Simulator/Core/GPUModel.cpp
Original file line number Diff 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 @@ -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;
}
4 changes: 4 additions & 0 deletions Simulator/Core/GPUModel.h
Original file line number Diff 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
34 changes: 34 additions & 0 deletions Simulator/Core/OperationManager.cpp
Original file line number Diff 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 @@ -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
}
}
}
Expand All @@ -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()";
}
Expand Down
7 changes: 7 additions & 0 deletions Simulator/Core/OperationManager.h
Original file line number Diff 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
3 changes: 2 additions & 1 deletion Simulator/Core/Operations.h
Original file line number Diff line number Diff line change
Expand Up @@ -22,5 +22,6 @@ enum class Operations {
copyToGPU,
copyFromGPU,
allocateGPU,
registerHistoryVariables
registerHistoryVariables,
loadEpochInputs
};
Loading
Loading