Implement getStepCounter in Host and Device APIs (#272)

* Host api access to stepCounter

Requires changing where step_count is incremented to the end of step

* Add host_api tests for getStepCounter

Covers init, host, step, exit functions and exit conditions

* Dont get curve instances to use a static method

* Remove non-useful test which caused errors

* Implement getStepCounter in device code

With AgentFunction, AgentCondition and RTCAgentFunction tests

Modifies RTC environment usage to be based on the existing EnvironmentProperties from the CUDAAgentModel rather than the EnvironmentDescription alone.

include/flamegpu/gpu/CUDAAgentModel.h

@@ -165,6 +165,10 @@ class CUDAAgentModel : public Simulation {
      * Number of times step() has been called since sim was last reset/init
      */
     unsigned int step_count;
+    /**
+     * Update the step counter for host and device.
+     */
+    void incrementStepCounter();
     /**
      * Map of agent storage 
      */

include/flamegpu/runtime/flamegpu_device_api.h

@@ -65,6 +65,14 @@ class FLAMEGPU_READ_ONLY_DEVICE_API {
      */
     const DeviceEnvironment environment;
 
+    /**
+     * Access the current stepCount
+     * @return the current step count, 0 indexed unsigned.
+     */
+    __forceinline__ __device__ unsigned int getStepCounter() const {
+        return environment.get<unsigned int>("_stepCount");
+    }
+
  protected:
     Curve::NamespaceHash agent_func_name_hash;
 

include/flamegpu/runtime/flamegpu_host_api.h

@@ -75,6 +75,12 @@ class FLAMEGPU_HOST_API {
      */
     const HostEnvironment environment;
 
+    /**
+     * Access the current stepCount
+     * @return the current step count, 0 indexed unsigned.
+     */
+    unsigned int getStepCounter() const;
+
  private:
     /**
      * Used internally for tracking what CUB has already calculated temp memory for

include/flamegpu/runtime/utility/DeviceEnvironment.cuh

@@ -93,7 +93,9 @@ __device__ __forceinline__ T DeviceEnvironment::get(const char(&name)[N]) const
     } else {
         curve_internal::d_curve_error = Curve::DEVICE_ERROR_UNKNOWN_VARIABLE;
         assert(false);
-        return *reinterpret_cast<T*>(reinterpret_cast<void*>(&flamegpu_internal::c_deviceEnvErrorPattern));
+        T rtn;
+        memcpy(&rtn, &flamegpu_internal::c_deviceEnvErrorPattern, sizeof(T));
+        return rtn;
     }
 }
 template<typename T, unsigned int N>
@@ -108,7 +110,9 @@ __device__ __forceinline__ T DeviceEnvironment::get(const char(&name)[N], const
     } else {
         curve_internal::d_curve_error = Curve::DEVICE_ERROR_UNKNOWN_VARIABLE;
         assert(false);
-        return *reinterpret_cast<T*>(reinterpret_cast<void*>(&flamegpu_internal::c_deviceEnvErrorPattern));
+        T rtn;
+        memcpy(&rtn, &flamegpu_internal::c_deviceEnvErrorPattern, sizeof(T));
+        return rtn;
     }
 }
 

include/flamegpu/runtime/utility/EnvironmentManager.cuh

@@ -22,6 +22,7 @@
 
 class EnvironmentDescription;
 class CUDAAgentModel;
+class CUDAAgent;
 
 
 /**
@@ -33,6 +34,10 @@ class CUDAAgentModel;
  * @note Not thread-safe
  */
 class EnvironmentManager {
+    /**
+     * Uses instance to for RTC compilation
+     */
+    friend class CUDAAgent;
     /**
      * Uses instance to initialise a models environment properties on the device
      */
@@ -148,8 +153,10 @@ class EnvironmentManager {
     /**
      * RTC functions hold thier own unique constants for environment variables. This function copies all environment variable to the RTC copies.
      * It can not be incorporated into init() as init will be called before RTC functions have been compiled.
+     * Uses the already populated Environment data from the cuda_model rather than environmentDescription.
+     * @param cuda_model the cuda model being initialised.
      */
-    void initRTC(const CUDAAgentModel& cuda_model, const EnvironmentDescription& desc);
+    void initRTC(const CUDAAgentModel& cuda_model);
     /**
      * Deactives all environmental properties linked to the named model from constant cache
      * @param model_name Name of the model

src/flamegpu/gpu/CUDAAgent.cu

@@ -219,16 +219,16 @@ void CUDAAgent::mapRuntimeVariables(const AgentFunctionData& func, const std::st
             agent_description.name.c_str(), func.initial_state.c_str());
     }
 
-    const Curve::VariableHash agent_hash = Curve::getInstance().variableRuntimeHash(agent_description.name.c_str());
-    const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str());
+    const Curve::VariableHash agent_hash = Curve::variableRuntimeHash(agent_description.name.c_str());
+    const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str());
     const unsigned int agent_count = this->getStateSize(func.initial_state);
     // loop through the agents variables to map each variable name using cuRVE
     for (const auto &mmp : agent_description.variables) {
         // get a device pointer for the agent variable name
         void* d_ptr = sm->second->getAgentListVariablePointer(mmp.first);
 
         // map using curve
-        const Curve::VariableHash var_hash = Curve::getInstance().variableRuntimeHash(mmp.first.c_str());
+        const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str());
 
         // get the agent variable size
         const size_t type_size = mmp.second.type_size * mmp.second.elements;
@@ -271,15 +271,15 @@ void CUDAAgent::unmapRuntimeVariables(const AgentFunctionData& func) const {
             agent_description.name.c_str(), func.initial_state.c_str());
     }
 
-    const Curve::VariableHash agent_hash = Curve::getInstance().variableRuntimeHash(agent_description.name.c_str());
-    const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str());
+    const Curve::VariableHash agent_hash = Curve::variableRuntimeHash(agent_description.name.c_str());
+    const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str());
     // loop through the agents variables to map each variable name using cuRVE
     for (const auto &mmp : agent_description.variables) {
         // get a device pointer for the agent variable name
         // void* d_ptr = sm->second->getAgentListVariablePointer(mmp.first);
 
         // unmap using curve
-        const Curve::VariableHash var_hash = Curve::getInstance().variableRuntimeHash(mmp.first.c_str());
+        const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str());
         Curve::getInstance().unregisterVariableByHash(var_hash + agent_hash + func_hash);
     }
 
@@ -434,15 +434,15 @@ void CUDAAgent::mapNewRuntimeVariables(const AgentFunctionData& func, const unsi
             agent_description.name.c_str(), func.agent_output_state.c_str());
     }
 
-    const Curve::VariableHash _agent_birth_hash = Curve::getInstance().variableRuntimeHash("_agent_birth");
-    const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str());
+    const Curve::VariableHash _agent_birth_hash = Curve::variableRuntimeHash("_agent_birth");
+    const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str());
     // loop through the agents variables to map each variable name using cuRVE
     for (const auto &mmp : agent_description.variables) {
         // get a device pointer for the agent variable name
         void* d_ptr = sm->second->getAgentNewListVariablePointer(mmp.first);
 
         // map using curve
-        const Curve::VariableHash var_hash = Curve::getInstance().variableRuntimeHash(mmp.first.c_str());
+        const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str());
 
         // get the agent variable size
         size_t type_size = mmp.second.type_size * mmp.second.elements;
@@ -464,15 +464,15 @@ void CUDAAgent::mapNewRuntimeVariables(const AgentFunctionData& func, const unsi
 }
 
 void CUDAAgent::unmapNewRuntimeVariables(const AgentFunctionData& func) const {
-    const Curve::VariableHash _agent_birth_hash = Curve::getInstance().variableRuntimeHash("_agent_birth");
-    const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str());
+    const Curve::VariableHash _agent_birth_hash = Curve::variableRuntimeHash("_agent_birth");
+    const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str());
     // loop through the agents variables to map each variable name using cuRVE
     for (const auto &mmp : agent_description.variables) {
         // get a device pointer for the agent variable name
         // void* d_ptr = sm->second->getAgentListVariablePointer(mmp.first);
 
         // unmap using curve
-        const Curve::VariableHash var_hash = Curve::getInstance().variableRuntimeHash(mmp.first.c_str());
+        const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str());
         Curve::getInstance().unregisterVariableByHash(var_hash + _agent_birth_hash + func_hash);
 
         // no need to unmap RTC variables
@@ -575,8 +575,8 @@ void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool f
     // curve rtc header
     CurveRTCHost curve_header;
     // agent function hash
-    Curve::NamespaceHash agentname_hash = Curve::getInstance().variableRuntimeHash(this->getAgentDescription().name.c_str());
-    Curve::NamespaceHash funcname_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str());
+    Curve::NamespaceHash agentname_hash = Curve::variableRuntimeHash(this->getAgentDescription().name.c_str());
+    Curve::NamespaceHash funcname_hash = Curve::variableRuntimeHash(func.name.c_str());
     Curve::NamespaceHash agent_func_name_hash = agentname_hash + funcname_hash;
 
     // set agent function variables in rtc curve
@@ -589,7 +589,7 @@ void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool f
         // Set input message variables in curve
         if (auto im = func.message_input.lock()) {
             // get the message input hash
-            Curve::NamespaceHash msg_in_hash = Curve::getInstance().variableRuntimeHash(im->name.c_str());
+            Curve::NamespaceHash msg_in_hash = Curve::variableRuntimeHash(im->name.c_str());
             for (auto msg_in_var : im->variables) {
                 // register message variables using combined hash
                 curve_header.registerVariable(msg_in_var.first.c_str(), msg_in_hash + agent_func_name_hash, msg_in_var.second.type.name(), msg_in_var.second.elements, true, false);
@@ -598,7 +598,7 @@ void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool f
         // Set output message variables in curve
         if (auto om = func.message_output.lock()) {
             // get the message input hash
-            Curve::NamespaceHash msg_out_hash = Curve::getInstance().variableRuntimeHash(om->name.c_str());
+            Curve::NamespaceHash msg_out_hash = Curve::variableRuntimeHash(om->name.c_str());
             for (auto msg_out_var : om->variables) {
                 // register message variables using combined hash
                 curve_header.registerVariable(msg_out_var.first.c_str(), msg_out_hash + agent_func_name_hash, msg_out_var.second.type.name(), msg_out_var.second.elements, false, true);
@@ -607,17 +607,23 @@ void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool f
         // Set agent output variables in curve
         if (auto ao = func.agent_output.lock()) {
             // get the message input hash
-            Curve::NamespaceHash agent_out_hash = Curve::getInstance().variableRuntimeHash("_agent_birth");
+            Curve::NamespaceHash agent_out_hash = Curve::variableRuntimeHash("_agent_birth");
             for (auto agent_out_var : ao->variables) {
                 // register message variables using combined hash
                 curve_header.registerVariable(agent_out_var.first.c_str(), agent_out_hash + funcname_hash, agent_out_var.second.type.name(), agent_out_var.second.elements, false, true);
             }
         }
     }
+
     // Set Environment variables in curve
-    Curve::NamespaceHash model_hash = Curve::getInstance().variableRuntimeHash(cuda_model.getModelDescription().name.c_str());
-    for (auto prop : cuda_model.getModelDescription().environment->getPropertiesMap()) {
-        curve_header.registerEnvVariable(prop.first.c_str(), model_hash, prop.second.type.name(), prop.second.elements);
+    Curve::NamespaceHash model_hash = Curve::variableRuntimeHash(cuda_model.getModelDescription().name.c_str());
+    for (auto p : EnvironmentManager::getInstance().getPropertiesMap()) {
+        if (p.first.first == cuda_model.getModelDescription().name) {
+            const char* variableName = p.first.second.c_str();
+            const char* type = p.second.type.name();
+            unsigned int elements = p.second.elements;
+            curve_header.registerEnvVariable(variableName, model_hash, type, elements);
+        }
     }
 
     // get the dynamically generated header from curve rtc

src/flamegpu/gpu/CUDAAgentModel.cu

@@ -70,8 +70,6 @@ bool CUDAAgentModel::step() {
         fprintf(stdout, "Processing Simulation Step %u\n", step_count);
     }
 
-    step_count++;
-
     unsigned int nStreams = 1;
     std::string message_name;
     Curve::NamespaceHash message_name_inp_hash = 0;
@@ -482,10 +480,12 @@ bool CUDAAgentModel::step() {
 #ifdef VISUALISATION
             if (visualisation) {
                 NVTX_PUSH("CUDAAgentModel::step::ExitConditions::UpdateVisualisation");
-                visualisation->updateBuffers(step_count);
+                visualisation->updateBuffers(step_count+1);
                 NVTX_POP();
             }
 #endif
+            // If there were any exit conditions, we also need to update the step count
+            incrementStepCounter();
             return false;
         }
     // If we have exit conditions functions, we might have host agent creation
@@ -496,10 +496,12 @@ bool CUDAAgentModel::step() {
 #ifdef VISUALISATION
         if (visualisation) {
             NVTX_PUSH("CUDAAgentModel::step::UpdateVisualisation");
-            visualisation->updateBuffers(step_count);
+            visualisation->updateBuffers(step_count+1);
             NVTX_POP();
         }
 #endif
+    // Update step count at the end of the step - when it has completed.
+    incrementStepCounter();
     return true;
 }
 
@@ -774,14 +776,16 @@ void CUDAAgentModel::initialiseSingletons() {
 
         // Populate the environment properties in constant Cache
         singletons->environment.init(model->name, *model->environment);
+        // Add the CUDAAgentModel specific variables(s)
+        singletons->environment.add({model->name, "_stepCount"}, 0u, false);
 
         // Reinitialise random for this simulation instance
         singletons->rng.reseed(getSimulationConfig().random_seed);
 
         singletonsInitialised = true;
     }
 
-    // init RTC
+    // Ensure RTC is set up.
     initialiseRTC();
 }
 
@@ -809,7 +813,7 @@ void CUDAAgentModel::initialiseRTC() {
         }
 
         // Initialise device environment for RTC
-        singletons->environment.initRTC(*this, *model->environment);
+        singletons->environment.initRTC(*this);
 
         rtcInitialised = true;
     }
@@ -961,3 +965,7 @@ void CUDAAgentModel::RTCSetEnvironmentVariable(const char* variable_name, const
     }
 }
 
+void CUDAAgentModel::incrementStepCounter() {
+    this->step_count++;
+    this->singletons->environment.set({model->name, "_stepCount"}, this->step_count);
+}

src/flamegpu/runtime/flamegpu_host_api.cu

@@ -88,3 +88,13 @@ void FLAMEGPU_HOST_API::resizeTempStorage(const CUB_Config &cc, const unsigned i
     cub_largestAllocatedOp[cc] = items;
 }
 
+
+
+/**
+ * Access the current stepCount
+ * Sepearate implementation to avoid dependency loop with cuda agent model.
+ * @return the current step count, 0 indexed unsigned.
+ */
+unsigned int FLAMEGPU_HOST_API::getStepCounter() const {
+    return agentModel.getStepCounter();
+}

src/flamegpu/runtime/utility/EnvironmentManager.cu

@@ -71,7 +71,7 @@ void EnvironmentManager::init(const std::string& model_name, const EnvironmentDe
     defragment(&orderedProperties);
 }
 
-void EnvironmentManager::initRTC(const CUDAAgentModel& cuda_model, const EnvironmentDescription& desc) {
+void EnvironmentManager::initRTC(const CUDAAgentModel& cuda_model) {
     // check to ensure that model name is not already registered
     auto res = cuda_agent_models.find(cuda_model.getModelDescription().name);
     if (res != cuda_agent_models.end()) {
@@ -80,10 +80,15 @@ void EnvironmentManager::initRTC(const CUDAAgentModel& cuda_model, const Environ
     // register model name
     cuda_agent_models.emplace(cuda_model.getModelDescription().name, cuda_model);
 
-    // loop through environment properties
-    for (auto p : desc.getPropertiesMap()) {
-        // Register variable for use in any RTC functions
-        cuda_model.RTCSetEnvironmentVariable(p.first.c_str(), p.second.data.ptr , p.second.data.length);
+    // loop through environment properties, already registered by cuda_
+    for (auto &p : properties) {
+        if (p.first.first == cuda_model.getModelDescription().name) {
+            auto var_name = p.first.second;
+            auto src = hc_buffer + p.second.offset;
+            auto length = p.second.length;
+            // Register variable for use in any RTC functions
+            cuda_model.RTCSetEnvironmentVariable(var_name.c_str(), src, length);
+        }
     }
 }
 

tests/CMakeLists.txt

@@ -106,6 +106,7 @@ SET(TEST_CASE_SRC
     ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_agent_creation.cu
     ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_api.cu
     ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_environment_manager.cu
+    ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_api.cu
     ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_agent_creation.cu
     ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_environment.cu
     ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_random.cu