From 97a3fa307a190b19028e4fc1972153386e83b90f Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Tue, 26 Sep 2023 15:28:36 +0100 Subject: [PATCH 1/8] DeviceAPI::isAgent() DeviceAPI::isState() With new C/Python tests Closes #1108 Currently only RTC implemented with Python test. --- include/flamegpu/runtime/DeviceAPI.cuh | 40 +++++++++++++++++ .../runtime/detail/curve/DeviceCurve.cuh | 26 +++++++++++ .../runtime/detail/curve/curve_rtc.cuh | 18 ++++++++ .../runtime/detail/curve/curve_rtc.cpp | 28 ++++++++++++ src/flamegpu/simulation/detail/CUDAAgent.cu | 5 ++- tests/python/runtime/test_device_api.py | 45 ++++++++++++++++++- 6 files changed, 160 insertions(+), 2 deletions(-) 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/curve/DeviceCurve.cuh b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh index 5d23d96f3..20ee3d196 100644 --- a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh +++ b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh @@ -244,6 +244,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 +415,15 @@ template(name, Curve::variableHash("_macro_environment"), 0); } + +__device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { + return strcmp(agent_name, "todo") == 0; // @todo +} +__device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { + return strcmp(agent_state, "todo") == 0; // @todo +} + + } // 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/src/flamegpu/runtime/detail/curve/curve_rtc.cpp b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp index a24d7ee2d..424ff428d 100644 --- a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp +++ b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp @@ -106,6 +106,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 +172,22 @@ __device__ __forceinline__ void DeviceCurve::setNewAgentArrayVariable(const char $DYNAMIC_SETNEWAGENTARRAYVARIABLE_IMPL } +// https://stackoverflow.com/a/34873763/1646387 +__device__ __forceinline__ int strcmp(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); +} +__device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { + return strcmp(agent_name, "$DYNAMIC_AGENT_NAME") == 0; +} +__device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { + return strcmp(agent_state, "$DYNAMIC_AGENT_STATE") == 0; +} + } // namespace curve } // namespace detail } // namespace flamegpu @@ -331,6 +349,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 +948,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/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/tests/python/runtime/test_device_api.py b/tests/python/runtime/test_device_api.py index a165fe36f..b00925565 100644 --- a/tests/python/runtime/test_device_api.py +++ b/tests/python/runtime/test_device_api.py @@ -49,6 +49,17 @@ class DeviceAPITest(TestCase): return flamegpu::ALIVE; } """ + + + 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") @@ -139,6 +150,7 @@ def test_array_set(self): assert output_array[3] == 16 + j + def test_array_get(self): model = pyflamegpu.ModelDescription("test_array_get") agent = model.newAgent("agent_name") @@ -184,4 +196,35 @@ def test_array_get(self): assert instance.getVariableInt("a2") == 4 + j 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 From 7539dc2b8bbaef14dd8cab08755ec8dca9a1f57a Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Tue, 26 Sep 2023 15:42:10 +0100 Subject: [PATCH 2/8] add C test, it's a clone of Python test. --- .../runtime/detail/curve/DeviceCurve.cuh | 10 +++++ tests/python/runtime/test_device_api.py | 9 ++-- tests/test_cases/runtime/test_device_api.cu | 43 +++++++++++++++++++ 3 files changed, 56 insertions(+), 6 deletions(-) diff --git a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh index 20ee3d196..9a4954a09 100644 --- a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh +++ b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh @@ -416,6 +416,16 @@ __device__ __forceinline__ char* DeviceCurve::getEnvironmentMacroProperty(const return getVariablePtr(name, Curve::variableHash("_macro_environment"), 0); } +// https://stackoverflow.com/a/34873763/1646387 +__device__ __forceinline__ int strcmp(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); +} + __device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { return strcmp(agent_name, "todo") == 0; // @todo } diff --git a/tests/python/runtime/test_device_api.py b/tests/python/runtime/test_device_api.py index b00925565..47eebed21 100644 --- a/tests/python/runtime/test_device_api.py +++ b/tests/python/runtime/test_device_api.py @@ -49,8 +49,7 @@ class DeviceAPITest(TestCase): return flamegpu::ALIVE; } """ - - + 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"))); @@ -104,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): @@ -148,7 +146,6 @@ 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): @@ -196,8 +193,8 @@ def test_array_get(self): assert instance.getVariableInt("a2") == 4 + j 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") diff --git a/tests/test_cases/runtime/test_device_api.cu b/tests/test_cases/runtime/test_device_api.cu index 5942ed2ad..c6015ddb3 100644 --- a/tests/test_cases/runtime/test_device_api.cu +++ b/tests/test_cases/runtime/test_device_api.cu @@ -362,5 +362,48 @@ 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 + const unsigned int agentCount = 100; + AgentVector init_population(agent, agentCount); + + // 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 From 6a2672d523f0b8ee9b3c467c962ac0bcf9baf762 Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Wed, 27 Sep 2023 10:09:35 +0100 Subject: [PATCH 3/8] Agent python Lint fix for prev C test. --- swig/python/codegen/codegen.py | 2 +- tests/python/codegen/test_codegen.py | 17 ++++++++++++++++- tests/test_cases/runtime/test_device_api.cu | 4 ++-- 3 files changed, 19 insertions(+), 4 deletions(-) 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/test_cases/runtime/test_device_api.cu b/tests/test_cases/runtime/test_device_api.cu index c6015ddb3..48edd552e 100644 --- a/tests/test_cases/runtime/test_device_api.cu +++ b/tests/test_cases/runtime/test_device_api.cu @@ -385,11 +385,11 @@ TEST(DeviceAPITest, check_agent_name_state) { // Init pop const unsigned int agentCount = 100; AgentVector init_population(agent, agentCount); - + // 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 From 5e15aeb90d6bd5296d448c59367d88ce4729428f Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Wed, 27 Sep 2023 10:28:22 +0100 Subject: [PATCH 4/8] Unify the implementation of dstrcmp() --- .../runtime/detail/curve/DeviceCurve.cuh | 16 +++---------- include/flamegpu/util/dstring.h | 24 +++++++++++++++++++ src/CMakeLists.txt | 1 + .../runtime/detail/curve/curve_rtc.cpp | 14 +++-------- 4 files changed, 31 insertions(+), 24 deletions(-) create mode 100644 include/flamegpu/util/dstring.h diff --git a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh index 9a4954a09..6fca20bc4 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__ @@ -416,24 +417,13 @@ __device__ __forceinline__ char* DeviceCurve::getEnvironmentMacroProperty(const return getVariablePtr(name, Curve::variableHash("_macro_environment"), 0); } -// https://stackoverflow.com/a/34873763/1646387 -__device__ __forceinline__ int strcmp(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); -} - __device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { - return strcmp(agent_name, "todo") == 0; // @todo + return dstrcmp(agent_name, "todo") == 0; // @todo } __device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { - return strcmp(agent_state, "todo") == 0; // @todo + return dstrcmp(agent_state, "todo") == 0; // @todo } - } // namespace curve } // namespace detail } // namespace flamegpu diff --git a/include/flamegpu/util/dstring.h b/include/flamegpu/util/dstring.h new file mode 100644 index 000000000..9428a5fdd --- /dev/null +++ b/include/flamegpu/util/dstring.h @@ -0,0 +1,24 @@ +#ifndef INCLUDE_FLAMEGPU_UTIL_DSTRING_H_ +#define INCLUDE_FLAMEGPU_UTIL_DSTRING_H_ + +#include + +/** + * Device implementations of required functionality + */ + + +/** + * strcmp() - Compare two strings, return 0 if equal, otherwise return suggests order + * + * @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); +} +#endif // INCLUDE_FLAMEGPU_UTIL_DSTRING_H_ diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 823840cfb..585f70bae 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -261,6 +261,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 diff --git a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp index 424ff428d..282de7908 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 { @@ -172,20 +173,11 @@ __device__ __forceinline__ void DeviceCurve::setNewAgentArrayVariable(const char $DYNAMIC_SETNEWAGENTARRAYVARIABLE_IMPL } -// https://stackoverflow.com/a/34873763/1646387 -__device__ __forceinline__ int strcmp(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); -} __device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { - return strcmp(agent_name, "$DYNAMIC_AGENT_NAME") == 0; + return dstrcmp(agent_name, "$DYNAMIC_AGENT_NAME") == 0; } __device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { - return strcmp(agent_state, "$DYNAMIC_AGENT_STATE") == 0; + return dstrcmp(agent_state, "$DYNAMIC_AGENT_STATE") == 0; } } // namespace curve From 1a7c2b5097dd642178e90dff9ba6e4bc06d4d41b Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Wed, 27 Sep 2023 11:35:25 +0100 Subject: [PATCH 5/8] Fix doc ci error --- include/flamegpu/util/dstring.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/include/flamegpu/util/dstring.h b/include/flamegpu/util/dstring.h index 9428a5fdd..9fd1472ab 100644 --- a/include/flamegpu/util/dstring.h +++ b/include/flamegpu/util/dstring.h @@ -4,13 +4,16 @@ #include /** - * Device implementations of required functionality + * 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) { From 86e82ec6abb4ccd208a6c0e1ada105409d81b4f8 Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Wed, 27 Sep 2023 13:13:31 +0100 Subject: [PATCH 6/8] C implementation and test fix --- include/flamegpu/runtime/AgentFunction.cuh | 6 +++ .../runtime/AgentFunctionCondition.cuh | 6 +++ include/flamegpu/runtime/detail/SharedBlock.h | 2 + .../runtime/detail/curve/DeviceCurve.cuh | 4 +- include/flamegpu/simulation/CUDASimulation.h | 5 +++ .../simulation/detail/DeviceStrings.h | 43 ++++++++++++++++++ src/CMakeLists.txt | 2 + src/flamegpu/simulation/CUDASimulation.cu | 14 +++++- .../simulation/detail/DeviceStrings.cu | 44 +++++++++++++++++++ tests/test_cases/runtime/test_device_api.cu | 3 +- 10 files changed, 124 insertions(+), 5 deletions(-) create mode 100644 include/flamegpu/simulation/detail/DeviceStrings.h create mode 100644 src/flamegpu/simulation/detail/DeviceStrings.cu diff --git a/include/flamegpu/runtime/AgentFunction.cuh b/include/flamegpu/runtime/AgentFunction.cuh index ca9e2271d..e12d303d8 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, @@ -60,6 +62,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 +81,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..3a7933cd4 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, @@ -44,6 +46,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 +60,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/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 6fca20bc4..a26d69dcf 100644 --- a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh +++ b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh @@ -418,10 +418,10 @@ __device__ __forceinline__ char* DeviceCurve::getEnvironmentMacroProperty(const } __device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { - return dstrcmp(agent_name, "todo") == 0; // @todo + return dstrcmp(agent_name, sm()->agent_name) == 0; } __device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { - return dstrcmp(agent_state, "todo") == 0; // @todo + return dstrcmp(agent_state, sm()->state_name) == 0; } } // 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/src/CMakeLists.txt b/src/CMakeLists.txt index 585f70bae..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 @@ -311,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/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/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/tests/test_cases/runtime/test_device_api.cu b/tests/test_cases/runtime/test_device_api.cu index 48edd552e..feece6d14 100644 --- a/tests/test_cases/runtime/test_device_api.cu +++ b/tests/test_cases/runtime/test_device_api.cu @@ -383,8 +383,7 @@ TEST(DeviceAPITest, check_agent_name_state) { AgentFunctionDescription func = agent.newFunction("some_function", check_agent_name_state_fn); model.newLayer().addAgentFunction(func); // Init pop - const unsigned int agentCount = 100; - AgentVector init_population(agent, agentCount); + AgentVector init_population(agent, AGENT_COUNT); // Setup Model CUDASimulation cudaSimulation(model); From 83234ae90630ae37f09a0cab27155522ff15da02 Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Wed, 27 Sep 2023 13:24:20 +0100 Subject: [PATCH 7/8] Fix docs CI --- include/flamegpu/runtime/AgentFunction.cuh | 2 ++ include/flamegpu/runtime/AgentFunctionCondition.cuh | 2 ++ 2 files changed, 4 insertions(+) diff --git a/include/flamegpu/runtime/AgentFunction.cuh b/include/flamegpu/runtime/AgentFunction.cuh index e12d303d8..4f6d74095 100644 --- a/include/flamegpu/runtime/AgentFunction.cuh +++ b/include/flamegpu/runtime/AgentFunction.cuh @@ -42,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) diff --git a/include/flamegpu/runtime/AgentFunctionCondition.cuh b/include/flamegpu/runtime/AgentFunctionCondition.cuh index 3a7933cd4..c2287d3f7 100644 --- a/include/flamegpu/runtime/AgentFunctionCondition.cuh +++ b/include/flamegpu/runtime/AgentFunctionCondition.cuh @@ -32,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 From d58bef0477b1ea8786185e008e08c49ef4fd1cfe Mon Sep 17 00:00:00 2001 From: Robert Chisholm Date: Wed, 27 Sep 2023 13:56:09 +0100 Subject: [PATCH 8/8] Place dstring.h inside namespace. --- include/flamegpu/runtime/detail/curve/DeviceCurve.cuh | 4 ++-- include/flamegpu/util/dstring.h | 6 ++++++ src/flamegpu/runtime/detail/curve/curve_rtc.cpp | 4 ++-- 3 files changed, 10 insertions(+), 4 deletions(-) diff --git a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh index a26d69dcf..38e7258b8 100644 --- a/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh +++ b/include/flamegpu/runtime/detail/curve/DeviceCurve.cuh @@ -418,10 +418,10 @@ __device__ __forceinline__ char* DeviceCurve::getEnvironmentMacroProperty(const } __device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { - return dstrcmp(agent_name, sm()->agent_name) == 0; + return util::dstrcmp(agent_name, sm()->agent_name) == 0; } __device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { - return dstrcmp(agent_state, sm()->state_name) == 0; + return util::dstrcmp(agent_state, sm()->state_name) == 0; } } // namespace curve diff --git a/include/flamegpu/util/dstring.h b/include/flamegpu/util/dstring.h index 9fd1472ab..35b8e01b1 100644 --- a/include/flamegpu/util/dstring.h +++ b/include/flamegpu/util/dstring.h @@ -3,6 +3,8 @@ #include +namespace flamegpu { +namespace util { /** * Device implementations of required string.h functionality */ @@ -24,4 +26,8 @@ __device__ __forceinline__ int dstrcmp(const char *s1, const char *s2) { return (*p1 > *p2) - (*p2 > *p1); } + +} // namespace util +} // namespace flamegpu + #endif // INCLUDE_FLAMEGPU_UTIL_DSTRING_H_ diff --git a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp index 282de7908..f9cf4dc38 100644 --- a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp +++ b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp @@ -174,10 +174,10 @@ __device__ __forceinline__ void DeviceCurve::setNewAgentArrayVariable(const char } __device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) { - return dstrcmp(agent_name, "$DYNAMIC_AGENT_NAME") == 0; + return util::dstrcmp(agent_name, "$DYNAMIC_AGENT_NAME") == 0; } __device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) { - return dstrcmp(agent_state, "$DYNAMIC_AGENT_STATE") == 0; + return util::dstrcmp(agent_state, "$DYNAMIC_AGENT_STATE") == 0; } } // namespace curve