diff --git a/include/flamegpu/runtime/AgentFunction.cuh b/include/flamegpu/runtime/AgentFunction.cuh index ca9e2271d..4f6d74095 100644 --- a/include/flamegpu/runtime/AgentFunction.cuh +++ b/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 } diff --git a/include/flamegpu/runtime/AgentFunctionCondition.cuh b/include/flamegpu/runtime/AgentFunctionCondition.cuh index 88c3c97df..c2287d3f7 100644 --- a/include/flamegpu/runtime/AgentFunctionCondition.cuh +++ b/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 } diff --git a/include/flamegpu/runtime/DeviceAPI.cuh b/include/flamegpu/runtime/DeviceAPI.cuh index f7782f8b2..4f60d54c1 100644 --- a/include/flamegpu/runtime/DeviceAPI.cuh +++ b/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 diff --git a/include/flamegpu/runtime/detail/SharedBlock.h b/include/flamegpu/runtime/detail/SharedBlock.h index b1ba3c15d..b75a36612 100644 --- a/include/flamegpu/runtime/detail/SharedBlock.h +++ b/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; diff --git a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh index 5d23d96f3..38e7258b8 100644 --- a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh +++ b/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 __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(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 diff --git a/include/flamegpu/runtime/detail/curve/curve_rtc.cuh b/include/flamegpu/runtime/detail/curve/curve_rtc.cuh index 9ed71aea3..557bb753e 100644 --- a/include/flamegpu/runtime/detail/curve/curve_rtc.cuh +++ b/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 { * */ std::map RTCEnvMacroProperties; + /** + * Agent name for ReadOnlyDeviceAPI::isAgent() + */ + std::string agentName; + /** + * Agent name for ReadOnlyDeviceAPI::isState() + */ + std::string agentState; }; } // namespace curve diff --git a/include/flamegpu/simulation/CUDASimulation.h b/include/flamegpu/simulation/CUDASimulation.h index 2b090d6a5..edc03f6fb 100644 --- a/include/flamegpu/simulation/CUDASimulation.h +++ b/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 &environment) : environment(environment) { } } * singletons; diff --git a/include/flamegpu/simulation/detail/DeviceStrings.h b/include/flamegpu/simulation/detail/DeviceStrings.h new file mode 100644 index 000000000..360073268 --- /dev/null +++ b/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 +#include +#include + +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 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_ diff --git a/include/flamegpu/util/dstring.h b/include/flamegpu/util/dstring.h new file mode 100644 index 000000000..35b8e01b1 --- /dev/null +++ b/include/flamegpu/util/dstring.h @@ -0,0 +1,33 @@ +#ifndef INCLUDE_FLAMEGPU_UTIL_DSTRING_H_ +#define INCLUDE_FLAMEGPU_UTIL_DSTRING_H_ + +#include + +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_ diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 823840cfb..258b6e947 100644 --- a/src/CMakeLists.txt +++ b/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 diff --git a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp index a24d7ee2d..f9cf4dc38 100644 --- a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp +++ b/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 __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 @@ -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) { diff --git a/src/flamegpu/simulation/CUDASimulation.cu b/src/flamegpu/simulation/CUDASimulation.cu index ed06c02aa..202d2ec59 100644 --- a/src/flamegpu/simulation/CUDASimulation.cu +++ b/src/flamegpu/simulation/CUDASimulation.cu @@ -726,6 +726,8 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un error_buffer, #endif cuda_agent.getCurve(func_des->name + "_condition").getDevicePtr(), + this->singletons->strings.getDeviceString(func_agent->name), + this->singletons->strings.getDeviceString(func_des->initial_state), static_cast(this->singletons->environment->getDeviceBuffer()), state_list_size, t_rng, @@ -950,7 +952,9 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un error_buffer, #endif cuda_agent.getCurve(func_des->name).getDevicePtr(), - static_cast(this->singletons->environment->getDeviceBuffer()), + this->singletons->strings.getDeviceString(func_agent->name), + this->singletons->strings.getDeviceString(func_des->initial_state), + static_cast(this->singletons->environment->getDeviceBuffer()), d_agentOut_nextID, state_list_size, d_in_messagelist_metadata, @@ -1571,6 +1575,14 @@ void CUDASimulation::initialiseSingletons() { macro_env->init(*submodel->subenvironment, mastermodel->macro_env, stream_0); } + // Populate device strings + for (const auto &[agent_name, agent] : model->agents) { + singletons->strings.registerDeviceString(agent_name); + for (const auto &state_name : agent->states) { + singletons->strings.registerDeviceString(state_name); + } + } + // Propagate singleton init to submodels for (auto &sm : submodel_map) { sm.second->initialiseSingletons(); diff --git a/src/flamegpu/simulation/detail/CUDAAgent.cu b/src/flamegpu/simulation/detail/CUDAAgent.cu index 81ae9dbfa..a7f228d6d 100644 --- a/src/flamegpu/simulation/detail/CUDAAgent.cu +++ b/src/flamegpu/simulation/detail/CUDAAgent.cu @@ -132,7 +132,7 @@ void CUDAAgent::setPopulationData(const AgentVector& population, const std::stri if (state_name == ModelData::DEFAULT_STATE) { THROW exception::InvalidAgentState("Agent '%s' does not use the default state, so the state must be passed explicitly, " "in CUDAAgent::setPopulationData()", - state_name.c_str(), population.getAgentName().c_str()); + population.getAgentName().c_str()); } else { THROW exception::InvalidAgentState("State '%s' was not found in agent '%s', " "in CUDAAgent::setPopulationData()", @@ -502,6 +502,9 @@ void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, const // Set Environment macro properties in curve macro_env->mapRTCVariables(curve_header); + // Set the agent name/state + curve_header.registerAgent(this->agent_description.name, func.initial_state); + std::string header_filename = std::string(func.rtc_func_name).append("_impl"); if (function_condition) header_filename.append("_condition"); diff --git a/src/flamegpu/simulation/detail/DeviceStrings.cu b/src/flamegpu/simulation/detail/DeviceStrings.cu new file mode 100644 index 000000000..6abfff612 --- /dev/null +++ b/src/flamegpu/simulation/detail/DeviceStrings.cu @@ -0,0 +1,44 @@ +#include "flamegpu/simulation/detail/DeviceStrings.h" + +#include "flamegpu/detail/cuda.cuh" +#include "flamegpu/simulation/detail/CUDAErrorChecking.cuh" + +namespace flamegpu { +namespace detail { + +DeviceStrings::~DeviceStrings() { + gpuErrchk(detail::cuda::cudaFree(device_buffer)); +} +void DeviceStrings::registerDeviceString(const std::string &host_string) { + if (offsets.find(host_string) == offsets.end()) { + offsets.emplace(host_string, host_buffer.size()); + host_stream << host_string; + host_stream << '\0'; // Each string requires a null terminating char + host_buffer = host_stream.str(); + } +} +const char* DeviceStrings::getDeviceString(const std::string &host_string) { + if (offsets.find(host_string) == offsets.end()) { + registerDeviceString(host_string); + } + const size_t host_buffer_len = host_buffer.size(); + const ptrdiff_t device_string_offset = offsets.at(host_string); + // Reallocate device buffer if necessary + if (!device_buffer || device_buffer_len < host_buffer_len) { + // Double buffer len in size + device_buffer_len = device_buffer_len == 0 ? 1024 : device_buffer_len * 2; + gpuErrchk(cudaFree(device_buffer)); + gpuErrchk(cudaMalloc(&device_buffer, device_buffer_len)); + device_buffer_occupied = 0; + } + // Update device buffer if necessary + if (device_buffer_occupied < host_buffer_len) { + gpuErrchk(cudaMemcpy(device_buffer, host_buffer.c_str(), host_buffer_len, cudaMemcpyHostToDevice)); + device_buffer_occupied = host_buffer_len; + } + // Return + return device_buffer + device_string_offset; +} + +} // namespace detail +} // namespace flamegpu diff --git a/swig/python/codegen/codegen.py b/swig/python/codegen/codegen.py index 1bc014f1b..264c36314 100644 --- a/swig/python/codegen/codegen.py +++ b/swig/python/codegen/codegen.py @@ -77,7 +77,7 @@ class CodeGenerator: } # getVariableType and setVariableType functions are added dynamically - fgpu_funcs = [ "getID", "getStepCounter", "getIndex" ] + fgpu_funcs = [ "getID", "getStepCounter", "getIndex", "isAgent", "isState" ] fgpu_attrs = ["ALIVE", "DEAD"] fgpu_input_msg_funcs = ["radius", "at"] # functions that can be called on message_in that do NOT return iterators fgpu_input_msg_iter_funcs = ["wrap", "vn", "vn_wrap"] # functions that can be called on message_in that do return iterators diff --git a/tests/python/codegen/test_codegen.py b/tests/python/codegen/test_codegen.py index 7f6e4345f..d347be782 100644 --- a/tests/python/codegen/test_codegen.py +++ b/tests/python/codegen/test_codegen.py @@ -429,6 +429,19 @@ def func(message_in: pyflamegpu.MessageNone, message_out: pyflamegpu.MessageBrut } """ +py_fgpu_agent_func_check_agent_name_state = """\ +@pyflamegpu.agent_function +def func(message_in: pyflamegpu.MessageNone, message_out: pyflamegpu.MessageNone) : + a = pyflamegpu.isAgent("foo") + b = pyflamegpu.isState("bar"); +""" +cpp_fgpu_agent_func_check_agent_name_state = """\ +FLAMEGPU_AGENT_FUNCTION(func, flamegpu::MessageNone, flamegpu::MessageNone){ + auto a = FLAMEGPU->isAgent("foo"); + auto b = FLAMEGPU->isState("bar"); +} +""" + py_fgpu_device_func_args = """\ @pyflamegpu.device_function def func(x: int) -> int : @@ -871,8 +884,10 @@ def test_fgpu_agent_func_return_type(self): """ Return type on an agent function raises a warning not error """ self._checkWarning(py_fgpu_agent_func_return_type, cpp_fgpu_agent_func_return_type, "Function definition return type not supported") - # device functions, arg types and calling + def test_fgpu_agent_func_check_agent_name_state(self): + self._checkExpected(py_fgpu_agent_func_check_agent_name_state, cpp_fgpu_agent_func_check_agent_name_state) + # device functions, arg types and calling def test_fgpu_agent_func_condition(self): # check correct format self._checkExpected(py_fgpu_cond_func, cpp_fgpu_cond_func) diff --git a/tests/python/runtime/test_device_api.py b/tests/python/runtime/test_device_api.py index a165fe36f..47eebed21 100644 --- a/tests/python/runtime/test_device_api.py +++ b/tests/python/runtime/test_device_api.py @@ -50,6 +50,16 @@ class DeviceAPITest(TestCase): } """ + agent_fn_check_agent_name_state = """ + FLAMEGPU_AGENT_FUNCTION(check_agent_name_state, flamegpu::MessageNone, flamegpu::MessageNone){ + FLAMEGPU->setVariable("correct_name", static_cast(FLAMEGPU->isAgent("agent"))); + FLAMEGPU->setVariable("wrong_name", static_cast(FLAMEGPU->isAgent("agent3"))); + FLAMEGPU->setVariable("correct_state", static_cast(FLAMEGPU->isState("state"))); + FLAMEGPU->setVariable("wrong_state", static_cast(FLAMEGPU->isState("state5"))); + return flamegpu::ALIVE; + } + """ + def test_agent_death_array(self): model = pyflamegpu.ModelDescription("test_agent_death_array") agent = model.newAgent("agent_name") @@ -93,7 +103,6 @@ def test_agent_death_array(self): assert output_array[1] == 4 + j assert output_array[2] == 8 + j assert output_array[3] == 16 + j - def test_array_set(self): @@ -137,7 +146,7 @@ def test_array_set(self): assert output_array[1] == 4 + j assert output_array[2] == 8 + j assert output_array[3] == 16 + j - + def test_array_get(self): model = pyflamegpu.ModelDescription("test_array_get") @@ -185,3 +194,34 @@ def test_array_get(self): assert instance.getVariableInt("a3") == 8 + j assert instance.getVariableInt("a4") == 16 + j + + def test_check_agent_name_state(self): + model = pyflamegpu.ModelDescription("test_array_get") + agent = model.newAgent("agent") + agent.newState("state") + agent.newState("state8") + agent.newVariableInt("correct_name", -1) + agent.newVariableInt("wrong_name", -1) + agent.newVariableInt("correct_state", -1) + agent.newVariableInt("wrong_state", -1) + # Do nothing, but ensure variables are made available on device + func = agent.newRTCFunction("some_function", self.agent_fn_check_agent_name_state) + model.newLayer().addAgentFunction(func) + # Init pop + init_population = pyflamegpu.AgentVector(agent, AGENT_COUNT) + + # Setup Model + cudaSimulation = pyflamegpu.CUDASimulation(model) + cudaSimulation.setPopulationData(init_population, "state") + # Run 1 step to ensure data is pushed to device + cudaSimulation.step() + # Recover data from device + population = pyflamegpu.AgentVector(agent, AGENT_COUNT) + cudaSimulation.getPopulationData(population, "state") + # Check results are correct + assert len(population) == AGENT_COUNT + for instance in population: + assert instance.getVariableInt("correct_name") == 1 + assert instance.getVariableInt("wrong_name") == 0 + assert instance.getVariableInt("correct_state") == 1 + assert instance.getVariableInt("wrong_state") == 0 \ No newline at end of file diff --git a/tests/test_cases/runtime/test_device_api.cu b/tests/test_cases/runtime/test_device_api.cu index 5942ed2ad..feece6d14 100644 --- a/tests/test_cases/runtime/test_device_api.cu +++ b/tests/test_cases/runtime/test_device_api.cu @@ -362,5 +362,47 @@ TEST(DeviceAPITest, getStepCounterFunctionCondition) { } } +FLAMEGPU_AGENT_FUNCTION(check_agent_name_state_fn, MessageNone, MessageNone) { + FLAMEGPU->setVariable("correct_name", static_cast(FLAMEGPU->isAgent("agent"))); + FLAMEGPU->setVariable("wrong_name", static_cast(FLAMEGPU->isAgent("agent3"))); + FLAMEGPU->setVariable("correct_state", static_cast(FLAMEGPU->isState("state"))); + FLAMEGPU->setVariable("wrong_state", static_cast(FLAMEGPU->isState("state5"))); + return ALIVE; +} + +TEST(DeviceAPITest, check_agent_name_state) { + ModelDescription model("model"); + AgentDescription agent = model.newAgent("agent"); + agent.newState("state"); + agent.newState("state8"); + agent.newVariable("correct_name", -1); + agent.newVariable("wrong_name", -1); + agent.newVariable("correct_state", -1); + agent.newVariable("wrong_state", -1); + // Do nothing, but ensure variables are made available on device + AgentFunctionDescription func = agent.newFunction("some_function", check_agent_name_state_fn); + model.newLayer().addAgentFunction(func); + // Init pop + AgentVector init_population(agent, AGENT_COUNT); + + // Setup Model + CUDASimulation cudaSimulation(model); + cudaSimulation.setPopulationData(init_population, "state"); + + // Run 1 step to ensure data is pushed to device + cudaSimulation.step(); + // Recover data from device + AgentVector population(agent, AGENT_COUNT); + cudaSimulation.getPopulationData(population, "state"); + // Check results are correct + EXPECT_EQ(population.size(), AGENT_COUNT); + for (const auto &instance : population) { + EXPECT_EQ(instance.getVariable("correct_name"), 1); + EXPECT_EQ(instance.getVariable("wrong_name"), 0); + EXPECT_EQ(instance.getVariable("correct_state"), 1); + EXPECT_EQ(instance.getVariable("wrong_state"), 0); + } +} + } // namespace test_device_api } // namespace flamegpu