DeviceAPI::isAgent() DeviceAPI::isState()

include/flamegpu/runtime/AgentFunction.cuh

@@ -24,6 +24,8 @@ typedef void(AgentFunctionWrapper)(
 #endif
 #ifndef __CUDACC_RTC__
     const detail::curve::CurveTable *d_curve_table,
+    const char* d_agent_name,
+    const char* d_state_name,
     const char* d_env_buffer,
 #endif
     id_t *d_agent_output_nextID,
@@ -40,6 +42,8 @@ typedef void(AgentFunctionWrapper)(
  * Initialises FLAMEGPU_API instance
  * @param error_buffer Buffer used for detecting and reporting exception::DeviceErrors (flamegpu must be built with FLAMEGPU_SEATBELTS enabled for this to be used)
  * @param d_curve_table Pointer to curve hash table in device memory
+ * @param d_agent_name Pointer to agent name string
+ * @param d_state_name Pointer to agent state string
  * @param d_env_buffer Pointer to env buffer in device memory
  * @param d_agent_output_nextID If agent output is enabled, this points to a global memory src of the next suitable agent id, this will be atomically incremented at birth
  * @param popNo Total number of agents executing the function (number of threads launched)
@@ -60,6 +64,8 @@ __global__ void agent_function_wrapper(
 #endif
 #ifndef __CUDACC_RTC__
     const detail::curve::CurveTable* __restrict__ d_curve_table,
+    const char* d_agent_name,
+    const char* d_state_name,
     const char* d_env_buffer,
 #endif
     id_t *d_agent_output_nextID,
@@ -77,6 +83,8 @@ __global__ void agent_function_wrapper(
         sm()->device_exception = error_buffer;
 #endif
 #ifndef __CUDACC_RTC__
+        sm()->agent_name = d_agent_name;
+        sm()->state_name = d_state_name;
         sm()->env_buffer = d_env_buffer;
 #endif
     }

include/flamegpu/runtime/AgentFunctionCondition.cuh

@@ -19,6 +19,8 @@ typedef void(AgentFunctionConditionWrapper)(
 #endif
 #ifndef __CUDACC_RTC__
     const detail::curve::CurveTable* d_curve_table,
+    const char* d_agent_name,
+    const char* d_state_name,
     const char* d_env_buffer,
 #endif
     const unsigned int popNo,
@@ -30,6 +32,8 @@ typedef void(AgentFunctionConditionWrapper)(
  * Initialises FLAMEGPU_API instance
  * @param error_buffer Buffer used for detecting and reporting exception::DeviceErrors (flamegpu must be built with FLAMEGPU_SEATBELTS enabled for this to be used)
  * @param d_curve_table Pointer to curve hash table in device memory
+ * @param d_agent_name Pointer to agent name string
+ * @param d_state_name Pointer to agent state string
  * @param d_env_buffer Pointer to env buffer in device memory
  * @param popNo Total number of agents exeucting the function (number of threads launched)
  * @param d_rng Array of curand states for this kernel
@@ -44,6 +48,8 @@ __global__ void agent_function_condition_wrapper(
 #endif
 #ifndef __CUDACC_RTC__
     const detail::curve::CurveTable* __restrict__ d_curve_table,
+    const char* d_agent_name,
+    const char* d_state_name,
     const char* d_env_buffer,
 #endif
     const unsigned int popNo,
@@ -56,6 +62,8 @@ __global__ void agent_function_condition_wrapper(
         sm()->device_exception = error_buffer;
 #endif
 #ifndef __CUDACC_RTC__
+        sm()->agent_name = d_agent_name;
+        sm()->state_name = d_state_name;
         sm()->env_buffer = d_env_buffer;
 #endif
     }

include/flamegpu/runtime/DeviceAPI.cuh

@@ -129,6 +129,26 @@ class ReadOnlyDeviceAPI {
 #endif
         return blockIdx.x * blockDim.x + threadIdx.x;
     }
+    /**
+     * When passed an agent name, returns a boolean to confirm whether it matches the name of the current agent
+     *
+     * This function may be useful if an agent function is shared between multiple agents
+     *
+     * @note The performance of this function is unlikely to be cheap unless used as part of an RTC agent function.
+     */
+    __forceinline__ __device__ bool isAgent(const char* agent_name) {
+        return detail::curve::DeviceCurve::isAgent(agent_name);
+    }
+    /**
+     * When passed an agent state, returns a boolean to confirm whether it matches the name of the agent input state of the current agent function
+     *
+     * This function may be useful if an agent function is shared between multiple agent states
+     *
+     * @note The performance of this function is unlikely to be cheap unless used as part of an RTC agent function (whereby it can be processed at compile time).
+     */
+    __forceinline__ __device__ bool isState(const char* agent_state) {
+        return detail::curve::DeviceCurve::isState(agent_state);
+    }
 };
 
 /** @brief    A flame gpu api class for the device runtime only
@@ -336,6 +356,26 @@ class DeviceAPI {
 #endif
         return blockIdx.x * blockDim.x + threadIdx.x;
     }
+    /**
+     * When passed an agent name, returns a boolean to confirm whether it matches the name of the current agent
+     *
+     * This function may be useful if an agent function is shared between multiple agents
+     *
+     * @note The performance of this function is unlikely to be cheap unless used as part of an RTC agent function.
+     */
+    __forceinline__ __device__ bool isAgent(const char* agent_name) {
+        return detail::curve::DeviceCurve::isAgent(agent_name);
+    }
+    /**
+     * When passed an agent state, returns a boolean to confirm whether it matches the name of the agent input state of the current agent function
+     *
+     * This function may be useful if an agent function is shared between multiple agent states
+     *
+     * @note The performance of this function is unlikely to be cheap unless used as part of an RTC agent function (whereby it can be processed at compile time).
+     */
+    __forceinline__ __device__ bool isState(const char* agent_state) {
+        return detail::curve::DeviceCurve::isState(agent_state);
+    }
 
     /**
      * Provides access to message read functionality inside agent functions

include/flamegpu/runtime/detail/SharedBlock.h

@@ -22,6 +22,8 @@ struct SharedBlock {
     unsigned int curve_count[curve::Curve::MAX_VARIABLES];
 #endif
     const char* env_buffer;
+    const char* agent_name;
+    const char* state_name;
 #endif
 #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
     exception::DeviceExceptionBuffer *device_exception;

include/flamegpu/runtime/detail/curve/DeviceCurve.cuh

@@ -5,6 +5,7 @@
 #include "flamegpu/runtime/detail/curve/Curve.cuh"
 #include "flamegpu/exception/FLAMEGPUDeviceException_device.cuh"
 #include "flamegpu/detail/type_decode.h"
+#include "flamegpu/util/dstring.h"
 
 #ifdef FLAMEGPU_USE_GLM
 #ifdef __CUDACC__
@@ -244,6 +245,23 @@ class DeviceCurve {
      */
     template<typename T, unsigned int I = 1, unsigned int J = 1, unsigned int K = 1, unsigned int W = 1, unsigned int M>
     __device__ __forceinline__ static char *getEnvironmentMacroProperty(const char(&name)[M]);
+
+    /**
+     * When passed an agent name, returns a boolean to confirm whether it matches the name of the current agent
+     *
+     * This function may be useful if an agent function is shared between multiple agents
+     *
+     * @note The performance of this function is unlikely to be cheap unless used as part of an RTC agent function.
+     */
+    __device__ __forceinline__ static bool isAgent(const char* agent_name);
+    /**
+     * When passed an agent state, returns a boolean to confirm whether it matches the name of the agent input state of the current agent function
+     *
+     * This function may be useful if an agent function is shared between multiple agent states
+     *
+     * @note The performance of this function is unlikely to be cheap unless used as part of an RTC agent function (whereby it can be processed at compile time).
+     */
+    __device__ __forceinline__ static bool isState(const char* agent_state);
 };
 
 ////
@@ -398,6 +416,14 @@ template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned in
 __device__ __forceinline__ char* DeviceCurve::getEnvironmentMacroProperty(const char(&name)[M]) {
     return getVariablePtr<T, I*J*K*W>(name, Curve::variableHash("_macro_environment"), 0);
 }
+
+__device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) {
+    return util::dstrcmp(agent_name, sm()->agent_name) == 0;
+}
+__device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) {
+    return util::dstrcmp(agent_state, sm()->state_name) == 0;
+}
+
 }  // namespace curve
 }  // namespace detail
 }  // namespace flamegpu

include/flamegpu/runtime/detail/curve/curve_rtc.cuh

@@ -160,6 +160,16 @@ class CurveRTCHost {
      * @throws exception::UnknownInternalError If the specified property is not registered
      */
     void unregisterEnvMacroProperty(const char* propertyName);
+    /**
+     * Register the name of the agent and it's state of the agent function
+     *
+     * Used by ReadOnlyDeviceAPI::isAgent() and ReadOnlyDeviceAPI::isState()
+     *
+     * @param agentName Name of the agent
+     * @param agentState Name of the agent's state
+     * @throws exception::UnknownInternalError If the agent has already been registered
+     */
+    void registerAgent(const std::string &agentName, const std::string &agentState);
     /**
      * Set the filename tagged in the file (goes into a #line statement)
      * @param filename Name to be used for the file in compile errors
@@ -378,6 +388,14 @@ class CurveRTCHost {
      * <name, RTCVariableProperties>
      */
     std::map<std::string, RTCEnvMacroPropertyProperties> RTCEnvMacroProperties;
+    /**
+     * Agent name for ReadOnlyDeviceAPI::isAgent()
+     */
+    std::string agentName;
+    /**
+    * Agent name for ReadOnlyDeviceAPI::isState()
+    */
+    std::string agentState;
 };
 
 }  // namespace curve

include/flamegpu/simulation/CUDASimulation.h

@@ -20,6 +20,7 @@
 #include "flamegpu/runtime/agent/HostNewAgentAPI.h"
 #include "flamegpu/simulation/detail/CUDAMacroEnvironment.h"
 #include "flamegpu/simulation/detail/EnvironmentManager.cuh"
+#include "flamegpu/simulation/detail/DeviceStrings.h"
 
 #ifdef FLAMEGPU_VISUALISATION
 #include "flamegpu/visualiser/ModelVis.h"
@@ -566,6 +567,10 @@ class CUDASimulation : public Simulation {
          * Provides buffers for device error checking
          */
         exception::DeviceExceptionManager exception;
+        /**
+         * Provides copies of strings (agent/state names) on device
+         */
+        detail::DeviceStrings strings;
 #endif
         explicit Singletons(const std::shared_ptr<detail::EnvironmentManager> &environment) : environment(environment) { }
     } * singletons;

include/flamegpu/simulation/detail/DeviceStrings.h

@@ -0,0 +1,43 @@
+#ifndef INCLUDE_FLAMEGPU_SIMULATION_DETAIL_DEVICESTRINGS_H_
+#define INCLUDE_FLAMEGPU_SIMULATION_DETAIL_DEVICESTRINGS_H_
+
+#include <string>
+#include <map>
+#include <sstream>
+
+namespace flamegpu {
+namespace detail {
+/**
+ * Utility for copying strings to device
+ */
+class DeviceStrings {
+ public:
+    /**
+     * Deallocates held device pointers
+     */
+    ~DeviceStrings();
+    /**
+     * Register a device string
+     */
+    void registerDeviceString(const std::string &host_string);
+    /**
+     * Returns a device pointer to the provided string
+     * @note If reallocation is required, earlier pointers may be invalidated
+     */
+    const char* getDeviceString(const std::string &host_string);
+
+ private:
+    std::stringstream host_stream;
+    // Cache stream in a string to reduce stream->string repeat conversion during sim execution
+    std::string host_buffer;
+    // Hold the offset into buffer for all registered strings
+    std::map<std::string, ptrdiff_t> offsets;
+    char* device_buffer = nullptr;
+    size_t device_buffer_occupied = 0;
+    size_t device_buffer_len = 0;
+};
+
+}  // namespace detail
+}  // namespace flamegpu
+
+#endif  // INCLUDE_FLAMEGPU_SIMULATION_DETAIL_DEVICESTRINGS_H_

include/flamegpu/util/dstring.h

@@ -0,0 +1,33 @@
+#ifndef INCLUDE_FLAMEGPU_UTIL_DSTRING_H_
+#define INCLUDE_FLAMEGPU_UTIL_DSTRING_H_
+
+#include <cuda_runtime.h>
+
+namespace flamegpu {
+namespace util {
+/**
+ * Device implementations of required string.h functionality
+ */
+
+
+/**
+ * strcmp() - Compare two strings, return 0 if equal, otherwise return suggests order
+ *
+ * @param s1 First string to be compared
+ * @param s2 Second string to be compared
+ *
+ * @note Implementation based on https://stackoverflow.com/a/34873763/1646387
+ */
+__device__ __forceinline__ int dstrcmp(const char *s1, const char *s2) {
+    const unsigned char *p1 = (const unsigned char *)s1;
+    const unsigned char *p2 = (const unsigned char *)s2;
+
+    while (*p1 && *p1 == *p2) ++p1, ++p2;
+
+    return (*p1 > *p2) - (*p2  > *p1);
+}
+
+}  // namespace util
+}  // namespace flamegpu
+
+#endif  // INCLUDE_FLAMEGPU_UTIL_DSTRING_H_

src/CMakeLists.txt

@@ -206,6 +206,7 @@ SET(SRC_INCLUDE
     ${FLAMEGPU_ROOT}/include/flamegpu/simulation/RunPlan.h
     ${FLAMEGPU_ROOT}/include/flamegpu/simulation/RunPlanVector.h
     ${FLAMEGPU_ROOT}/include/flamegpu/simulation/Simulation.h
+    ${FLAMEGPU_ROOT}/include/flamegpu/simulation/detail/DeviceStrings.h
     ${FLAMEGPU_ROOT}/include/flamegpu/simulation/detail/EnvironmentManager.cuh
     ${FLAMEGPU_ROOT}/include/flamegpu/simulation/detail/RandomManager.cuh
     ${FLAMEGPU_ROOT}/include/flamegpu/runtime/AgentFunction.cuh
@@ -261,6 +262,7 @@ SET(SRC_INCLUDE
     ${FLAMEGPU_ROOT}/include/flamegpu/runtime/environment/HostEnvironment.cuh
     ${FLAMEGPU_ROOT}/include/flamegpu/runtime/environment/HostMacroProperty.cuh
     ${FLAMEGPU_ROOT}/include/flamegpu/util/cleanup.h
+    ${FLAMEGPU_ROOT}/include/flamegpu/util/dstring.h
     ${FLAMEGPU_ROOT}/include/flamegpu/util/nvtx.h
     ${FLAMEGPU_ROOT}/include/flamegpu/util/StringPair.h
     ${FLAMEGPU_ROOT}/include/flamegpu/detail/Any.h
@@ -310,6 +312,7 @@ SET(SRC_FLAMEGPU
     ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/CUDAMessage.cu
     ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/CUDAScatter.cu
     ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/CUDAMacroEnvironment.cu
+    ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/DeviceStrings.cu
     ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/SimRunner.cu
     ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/SimLogger.cu
     ${FLAMEGPU_ROOT}/src/flamegpu/simulation/detail/EnvironmentManager.cu

src/flamegpu/runtime/detail/curve/curve_rtc.cpp

@@ -28,6 +28,7 @@ const char* CurveRTCHost::curve_rtc_dynamic_h_template = R"###(dynamic/curve_rtc
 #include "flamegpu/exception/FLAMEGPUDeviceException.cuh"
 #include "flamegpu/detail/type_decode.h"
 #include "flamegpu/runtime/detail/curve/Curve.cuh"
+#include "flamegpu/util/dstring.h"
 
 namespace flamegpu {
 
@@ -106,6 +107,8 @@ class DeviceCurve {
     template <typename T, unsigned int N, unsigned int M>
     __device__ __forceinline__ static void setNewAgentArrayVariable(const char(&name)[M], T variable, unsigned int variable_index, unsigned int array_index);
 
+    __device__ __forceinline__ static bool isAgent(const char* agent_name);
+    __device__ __forceinline__ static bool isState(const char* agent_state);
 };
 
 template <typename T, unsigned int N>
@@ -170,6 +173,13 @@ __device__ __forceinline__ void DeviceCurve::setNewAgentArrayVariable(const char
 $DYNAMIC_SETNEWAGENTARRAYVARIABLE_IMPL    
 }
 
+__device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) {
+    return util::dstrcmp(agent_name, "$DYNAMIC_AGENT_NAME") == 0;
+}
+__device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) {
+    return util::dstrcmp(agent_state, "$DYNAMIC_AGENT_STATE") == 0;
+}
+
 }  // namespace curve 
 }  // namespace detail 
 }  // namespace flamegpu 
@@ -331,6 +341,14 @@ void CurveRTCHost::registerEnvVariable(const char* propertyName, ptrdiff_t offse
         THROW exception::UnknownInternalError("Environment property with name '%s' is already registered, in CurveRTCHost::registerEnvVariable()", propertyName);
     }
 }
+void CurveRTCHost::registerAgent(const std::string &_agentName, const std::string &_agentState) {
+    if (this->agentName.empty()) {
+        this->agentName = _agentName;
+        this->agentState = _agentState;
+    } else {
+        THROW exception::UnknownInternalError("Agent is already registered with name '%s' and state '%s', in CurveRTCHost::registerAgent()", this->agentName.c_str(), this->agentState.c_str());
+    }
+}
 
 void CurveRTCHost::unregisterEnvVariable(const char* propertyName) {
     auto i = RTCEnvVariables.find(propertyName);
@@ -922,6 +940,8 @@ void CurveRTCHost::initHeaderGetters() {
         getMessageArrayVariableLDGImpl << "           return {};\n";
         setHeaderPlaceholder("$DYNAMIC_GETMESSAGEARRAYVARIABLE_LDG_IMPL", getMessageArrayVariableLDGImpl.str());
     }
+    setHeaderPlaceholder("$DYNAMIC_AGENT_NAME", this->agentName);
+    setHeaderPlaceholder("$DYNAMIC_AGENT_STATE", this->agentState);
 }
 void CurveRTCHost::initDataBuffer() {
     if (data_buffer_size == 0 || h_data_buffer) {