diff --git a/include/flamegpu/io/JSONStateWriter.h b/include/flamegpu/io/JSONStateWriter.h index ba093c80f..f34c64db9 100644 --- a/include/flamegpu/io/JSONStateWriter.h +++ b/include/flamegpu/io/JSONStateWriter.h @@ -52,7 +52,7 @@ class JSONStateWriter : public StateWriter { void writeConfig(const Simulation *sim_instance) override; void writeStats(unsigned int iterations) override; void writeEnvironment(const std::shared_ptr& env_manager) override; - void writeMacroEnvironment(const std::shared_ptr& macro_env) override; + void writeMacroEnvironment(const std::shared_ptr& macro_env, std::initializer_list filter = {}) override; void writeAgents(const util::StringPairUnorderedMap>& agents_map) override; private: diff --git a/include/flamegpu/io/StateWriter.h b/include/flamegpu/io/StateWriter.h index d1c074454..f097c516b 100644 --- a/include/flamegpu/io/StateWriter.h +++ b/include/flamegpu/io/StateWriter.h @@ -81,7 +81,12 @@ class StateWriter { virtual void writeConfig(const Simulation *sim_instance) = 0; virtual void writeStats(unsigned int iterations) = 0; virtual void writeEnvironment(const std::shared_ptr& env_manager) = 0; - virtual void writeMacroEnvironment(const std::shared_ptr& macro_env) = 0; + /** + * Write the macro environment block + * @param macro_env The macro environment to pull properties from + * @param filter If provided, only named properties will be written. Note, if filter contains missing properties it will fail + */ + virtual void writeMacroEnvironment(const std::shared_ptr& macro_env, std::initializer_list filter = {}) = 0; virtual void writeAgents(const util::StringPairUnorderedMap>& agents_map) = 0; }; } // namespace io diff --git a/include/flamegpu/io/XMLStateWriter.h b/include/flamegpu/io/XMLStateWriter.h index 4ef7b0354..849b0c88a 100644 --- a/include/flamegpu/io/XMLStateWriter.h +++ b/include/flamegpu/io/XMLStateWriter.h @@ -46,7 +46,7 @@ class XMLStateWriter : public StateWriter { void writeConfig(const Simulation *sim_instance) override; void writeStats(unsigned int iterations) override; void writeEnvironment(const std::shared_ptr& env_manager) override; - void writeMacroEnvironment(const std::shared_ptr& macro_env) override; + void writeMacroEnvironment(const std::shared_ptr& macro_env, std::initializer_list filter = {}) override; void writeAgents(const util::StringPairUnorderedMap>& agents_map) override; private: diff --git a/include/flamegpu/runtime/environment/HostEnvironment.cuh b/include/flamegpu/runtime/environment/HostEnvironment.cuh index d09c82c33..93f6fd335 100644 --- a/include/flamegpu/runtime/environment/HostEnvironment.cuh +++ b/include/flamegpu/runtime/environment/HostEnvironment.cuh @@ -36,7 +36,9 @@ class HostEnvironment { /** * Constructor, to be called by HostAPI */ - explicit HostEnvironment(unsigned int instance_id, const std::shared_ptr &env, const std::shared_ptr& _macro_env); + explicit HostEnvironment(CUDASimulation &_simulation, cudaStream_t _stream, + std::shared_ptr env, + std::shared_ptr _macro_env); /** * Provides access to EnvironmentManager singleton */ @@ -50,6 +52,14 @@ class HostEnvironment { * This is used to augment all variable names */ const unsigned int instance_id; + /** + * The relevant simulation, required for importing macro properties + */ + CUDASimulation& simulation; + /** + * CUDAStream for memcpys + */ + const cudaStream_t stream; public: /** @@ -155,6 +165,21 @@ class HostEnvironment { template HostMacroProperty_swig getMacroProperty_swig(const std::string& name) const; #endif + /** + * Import macro property data from file + * @param property_name Name of the macro property to import + * @param file_path Path to file containing macro property data (.json, .xml, .bin) + * @note This method supports raw binary files (.bin) + */ + void importMacroProperty(const std::string& property_name, const std::string& file_path) const; + /** + * Export macro property data to file + * @param property_name Name of the macro property to import + * @param file_path Path to file to export macro property data (.json, .xml. bin) + * @param pretty_print Print in readable or minified format (if available) + * @note This method supports raw binary files (.bin) + */ + void exportMacroProperty(const std::string& property_name, const std::string& file_path, bool pretty_print = true) const; }; /** diff --git a/include/flamegpu/runtime/environment/HostMacroProperty.cuh b/include/flamegpu/runtime/environment/HostMacroProperty.cuh index 07cd5e8e3..1f0721deb 100644 --- a/include/flamegpu/runtime/environment/HostMacroProperty.cuh +++ b/include/flamegpu/runtime/environment/HostMacroProperty.cuh @@ -33,14 +33,22 @@ struct HostMacroProperty_MetaData { } /** * Download data + * @note Only works first time + * @see force_download() */ void download() { if (!h_base_ptr) { h_base_ptr = static_cast(malloc(elements * type_size)); - gpuErrchk(cudaMemcpyAsync(h_base_ptr, d_base_ptr, elements * type_size, cudaMemcpyDeviceToHost, stream)); - gpuErrchk(cudaStreamSynchronize(stream)); - has_changed = false; + force_download(); + } + } + void force_download() { + if (!h_base_ptr) { + h_base_ptr = static_cast(malloc(elements * type_size)); } + gpuErrchk(cudaMemcpyAsync(h_base_ptr, d_base_ptr, elements * type_size, cudaMemcpyDeviceToHost, stream)); + gpuErrchk(cudaStreamSynchronize(stream)); + has_changed = false; } /** * Upload data diff --git a/include/flamegpu/simulation/detail/CUDAMacroEnvironment.h b/include/flamegpu/simulation/detail/CUDAMacroEnvironment.h index 0f7cd7a41..98735436e 100644 --- a/include/flamegpu/simulation/detail/CUDAMacroEnvironment.h +++ b/include/flamegpu/simulation/detail/CUDAMacroEnvironment.h @@ -162,6 +162,13 @@ class CUDAMacroEnvironment { * Used for IO */ const std::map& getPropertiesMap() const; + /** + * Return the metadata, if setup, for the named macro property + * + * @param property_name Name of the macro property to load + * @return nullptr if not currently cached + */ + std::shared_ptr getHostPropertyMetadata(const std::string property_name); private: const CUDASimulation& cudaSimulation; diff --git a/src/flamegpu/io/JSONStateWriter.cu b/src/flamegpu/io/JSONStateWriter.cu index a138bf276..d4bfb71c6 100644 --- a/src/flamegpu/io/JSONStateWriter.cu +++ b/src/flamegpu/io/JSONStateWriter.cu @@ -208,7 +208,7 @@ void JSONStateWriter::writeEnvironment(const std::shared_ptr& macro_env) { +void JSONStateWriter::writeMacroEnvironment(const std::shared_ptr& macro_env, std::initializer_list filter) { if (!writer) { THROW exception::UnknownInternalError("beginWrite() must be called before writeMacroEnvironment(), in JSONStateWriter::writeMacroEnvironment()"); } else if (macro_environment_written) { @@ -220,6 +220,12 @@ void JSONStateWriter::writeMacroEnvironment(const std::shared_ptrStartObject(); if (macro_env) { const std::map& m_properties = macro_env->getPropertiesMap(); + for (const auto &_filter : filter) { + if (m_properties.find(_filter) == m_properties.end()) { + THROW exception::InvalidEnvProperty("Macro property '%s' specified in filter does not exist, in JSONStateWriter::writeMacroEnvironment()", _filter.c_str()); + } + } + std::set filter_set = filter; // Calculate largest buffer in map size_t max_len = 0; for (const auto& [_, prop] : m_properties) { @@ -230,6 +236,8 @@ void JSONStateWriter::writeMacroEnvironment(const std::shared_ptr(malloc(max_len)); // Write out each array (all are written out as 1D arrays for simplicity given variable dimensions) for (const auto& [name, prop] : m_properties) { + if (!filter_set.empty() && filter_set.find(name) == filter_set.end()) + continue; // Copy data const size_t element_ct = std::accumulate(prop.elements.begin(), prop.elements.end(), 1, std::multiplies()); gpuErrchk(cudaMemcpy(t_buffer, prop.d_ptr, element_ct * prop.type_size, cudaMemcpyDeviceToHost)); diff --git a/src/flamegpu/io/XMLStateWriter.cu b/src/flamegpu/io/XMLStateWriter.cu index 51e3b29aa..a73157202 100644 --- a/src/flamegpu/io/XMLStateWriter.cu +++ b/src/flamegpu/io/XMLStateWriter.cu @@ -259,7 +259,7 @@ void XMLStateWriter::writeEnvironment(const std::shared_ptr& macro_env) { +void XMLStateWriter::writeMacroEnvironment(const std::shared_ptr& macro_env, std::initializer_list filter) { if (!doc || !pRoot) { THROW exception::UnknownInternalError("beginWrite() must be called before writeMacroEnvironment(), in XMLStateWriter::writeMacroEnvironment()"); } else if (macro_environment_written) { @@ -269,6 +269,12 @@ void XMLStateWriter::writeMacroEnvironment(const std::shared_ptrNewElement("macro_environment"); if (macro_env) { const std::map& m_properties = macro_env->getPropertiesMap(); + for (const auto &_filter : filter) { + if (m_properties.find(_filter) == m_properties.end()) { + THROW exception::InvalidEnvProperty("Macro property '%s' specified in filter does not exist, in XMLStateWriter::writeMacroEnvironment()", _filter.c_str()); + } + } + std::set filter_set = filter; // Calculate largest buffer in map size_t max_len = 0; for (const auto& [_, prop] : m_properties) { @@ -279,6 +285,8 @@ void XMLStateWriter::writeMacroEnvironment(const std::shared_ptr(malloc(max_len)); // Write out each array (all are written out as 1D arrays for simplicity given variable dimensions) for (const auto& [name, prop] : m_properties) { + if (!filter_set.empty() && filter_set.find(name) == filter_set.end()) + continue; // Copy data const size_t element_ct = std::accumulate(prop.elements.begin(), prop.elements.end(), 1, std::multiplies()); gpuErrchk(cudaMemcpy(t_buffer, prop.d_ptr, element_ct * prop.type_size, cudaMemcpyDeviceToHost)); diff --git a/src/flamegpu/runtime/HostAPI.cu b/src/flamegpu/runtime/HostAPI.cu index 8b11fd1a7..57280e760 100644 --- a/src/flamegpu/runtime/HostAPI.cu +++ b/src/flamegpu/runtime/HostAPI.cu @@ -18,7 +18,7 @@ HostAPI::HostAPI(CUDASimulation &_agentModel, const unsigned int _streamId, cudaStream_t _stream) : random(rng) - , environment(_agentModel.getInstanceID(), env, macro_env) + , environment(_agentModel, _stream, env, macro_env) , agentModel(_agentModel) , d_output_space(nullptr) , d_output_space_size(0) diff --git a/src/flamegpu/runtime/environment/HostEnvironment.cu b/src/flamegpu/runtime/environment/HostEnvironment.cu index e63ce730e..39d9423fb 100644 --- a/src/flamegpu/runtime/environment/HostEnvironment.cu +++ b/src/flamegpu/runtime/environment/HostEnvironment.cu @@ -1,10 +1,113 @@ #include "flamegpu/runtime/environment/HostEnvironment.cuh" +#include +#include +#include +#include +#include + +#include "flamegpu/io/StateWriter.h" +#include "flamegpu/io/StateWriterFactory.h" +#include "flamegpu/io/StateReader.h" +#include "flamegpu/io/StateReaderFactory.h" +#include "flamegpu/simulation/CUDASimulation.h" + namespace flamegpu { -HostEnvironment::HostEnvironment(const unsigned int _instance_id, const std::shared_ptr &env, const std::shared_ptr& _macro_env) - : env_mgr(env) - , macro_env(_macro_env) - , instance_id(_instance_id) { } +HostEnvironment::HostEnvironment(CUDASimulation &_simulation, cudaStream_t _stream, + std::shared_ptr env, + std::shared_ptr _macro_env) + : env_mgr(std::move(env)) + , macro_env(std::move(_macro_env)) + , instance_id(_simulation.getInstanceID()) + , simulation(_simulation) + , stream(_stream) { } + +void HostEnvironment::importMacroProperty(const std::string& property_name, const std::string& file_path) const { + // Validate the property exists + const auto &m_props = macro_env->getPropertiesMap(); + const auto &m_prop = m_props.find(property_name); + if (m_prop == m_props.end()) { + THROW exception::InvalidEnvProperty("The environment macro property '%s' was not found within the model description, in HostEnvironment::importMacroProperty().", property_name.c_str()); + } + const unsigned int m_prop_elements = std::accumulate(m_prop->second.elements.begin(), m_prop->second.elements.end(), 1, std::multiplies()); + try { + io::StateReader *read__ = io::StateReaderFactory::createReader(file_path); + read__->parse(file_path, simulation.getModelDescription().shared_from_this(), Verbosity::Quiet); + std::unordered_map> macro_init; + read__->getMacroEnvironment(macro_init); + // Validate the property exists within macro_init + const auto &l_prop = macro_init.find(property_name); + if (l_prop == macro_init.end()) { + THROW exception::InvalidEnvProperty("The environment macro property '%s' was not found within the input file '%s'.", property_name.c_str(), file_path.c_str()); + } + // Check the length validates + if (l_prop->second.size() != m_prop_elements * m_prop->second.type_size) { + THROW exception::InvalidInputFile("Length of input file '%s's environment macro property '%s' does not match, (%u != %u), in HostEnvironment::importMacroProperty()", + file_path.c_str(), property_name.c_str(), static_cast(l_prop->second.size()), static_cast(m_prop_elements * m_prop->second.type_size)); + } + gpuErrchk(cudaMemcpyAsync(m_prop->second.d_ptr, l_prop->second.data(), l_prop->second.size(), cudaMemcpyHostToDevice, stream)); + } catch (const exception::UnsupportedFileType&) { + const std::string extension = std::filesystem::path(file_path).extension().string(); + if (extension == ".bin") { + // Additionally support raw binary dump + // Read the file + std::ifstream input(file_path, std::ios::binary); + std::vector buffer(std::istreambuf_iterator(input), {}); + // Check the length validates + if (buffer.size() != m_prop_elements * m_prop->second.type_size) { + THROW exception::InvalidInputFile("Length of binary input file '%s' does not match the environment macro property '%s', (%u != %u), in HostEnvironment::importMacroProperty()", + file_path.c_str(), property_name.c_str(), static_cast(buffer.size()), static_cast(m_prop_elements * m_prop->second.type_size)); + } + // Update the property + gpuErrchk(cudaMemcpyAsync(m_prop->second.d_ptr, buffer.data(), buffer.size(), cudaMemcpyHostToDevice, stream)); + } else { + throw; + } + } + gpuErrchk(cudaStreamSynchronize(stream)); + // If macro property exists in cache sync cache + if (const auto cache = macro_env->getHostPropertyMetadata(property_name)) { + cache->force_download(); + } +} +void HostEnvironment::exportMacroProperty(const std::string& property_name, const std::string& file_path, bool pretty_print) const { + // If macro property exists in cache sync cache + if (const auto cache = macro_env->getHostPropertyMetadata(property_name)) { + cache->upload(); + } + try { + io::StateWriter* write__ = io::StateWriterFactory::createWriter(file_path); + write__->beginWrite(file_path, pretty_print); + write__->writeMacroEnvironment(macro_env, { property_name }); + write__->endWrite(); + } catch (const exception::UnsupportedFileType&) { + const std::string extension = std::filesystem::path(file_path).extension().string(); + if (extension == ".bin") { + // Additionally support raw binary dump + // Validate the property exists + const auto &m_props = macro_env->getPropertiesMap(); + const auto &m_prop = m_props.find(property_name); + if (m_prop == m_props.end()) { + THROW exception::InvalidEnvProperty("The environment macro property '%s' was not found within the model description, in HostEnvironment::exportMacroProperty().", property_name.c_str()); + } + // Check the file doesn't already exist + if (std::filesystem::exists(file_path)) { + THROW exception::FileAlreadyExists("The binary output file '%s' already exists, in HostEnvironment::exportMacroProperty().", file_path.c_str()); + } + // Copy the data to a temporary buffer on host + const unsigned int m_prop_elements = std::accumulate(m_prop->second.elements.begin(), m_prop->second.elements.end(), 1, std::multiplies()); + std::vector buffer; + buffer.resize(m_prop_elements * m_prop->second.type_size); + gpuErrchk(cudaMemcpyAsync(buffer.data(), m_prop->second.d_ptr, m_prop_elements * m_prop->second.type_size, cudaMemcpyDeviceToHost, stream)); + gpuErrchk(cudaStreamSynchronize(stream)); + // Output to file + std::ofstream output(file_path, std::ios::binary); + output.write(buffer.data(), buffer.size()); + } else { + throw; + } + } +} } // namespace flamegpu diff --git a/src/flamegpu/simulation/detail/CUDAMacroEnvironment.cu b/src/flamegpu/simulation/detail/CUDAMacroEnvironment.cu index 49b9f8659..7606eaa26 100644 --- a/src/flamegpu/simulation/detail/CUDAMacroEnvironment.cu +++ b/src/flamegpu/simulation/detail/CUDAMacroEnvironment.cu @@ -102,6 +102,14 @@ void CUDAMacroEnvironment::unmapRTCVariables(detail::curve::CurveRTCHost& curve_ const std::map& CUDAMacroEnvironment::getPropertiesMap() const { return properties; } + +std::shared_ptr CUDAMacroEnvironment::getHostPropertyMetadata(const std::string property_name) { + auto cache = host_cache.find(property_name); + if (cache != host_cache.end()) { + return cache->second.lock(); + } + return nullptr; +} #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS void CUDAMacroEnvironment::resetFlagsAsync(const std::vector &streams) { unsigned int i = 0; diff --git a/tests/test_cases/io/test_io.cu b/tests/test_cases/io/test_io.cu index 8e8d833bf..7461f4e31 100644 --- a/tests/test_cases/io/test_io.cu +++ b/tests/test_cases/io/test_io.cu @@ -9,8 +9,11 @@ namespace flamegpu { namespace test_io { bool validate_has_run = false; +// Used by MiniSim3 where file name is required during a host function +const char* current_test_file_name = nullptr; const char *XML_FILE_NAME = "test.xml"; const char *JSON_FILE_NAME = "test.json"; +const char *BIN_FILE_NAME = "test.bin"; FLAMEGPU_STEP_FUNCTION(VALIDATE_ENV) { EXPECT_EQ(FLAMEGPU->environment.getProperty("float"), 12.0f); EXPECT_EQ(FLAMEGPU->environment.getProperty("float"), 12.0f); @@ -132,7 +135,32 @@ FLAMEGPU_STEP_FUNCTION(VALIDATE_MACRO_ENV) { auto macro_uint = FLAMEGPU->environment.getMacroProperty("macro_uint"); EXPECT_EQ(static_cast(macro_uint), 12u); } +FLAMEGPU_STEP_FUNCTION(VALIDATE_MACRO_ENV2) { + const auto macro_float_3_3_3 = FLAMEGPU->environment.getMacroProperty("macro_float_3_3_3"); + int i = 0; + for (int x = 0; x < 3; ++x) { + for (int y = 0; y < 3; ++y) { + for (int z = 0; z < 3; ++z) { + EXPECT_EQ(static_cast(macro_float_3_3_3[x][y][z]), static_cast(i++)); + } + } + } + auto macro_int_4_3_2_3 = FLAMEGPU->environment.getMacroProperty("macro_int_4_3_2_3"); + i = 0; + for (int x = 0; x < 4; ++x) { + for (int y = 0; y < 3; ++y) { + for (int z = 0; z < 2; ++z) { + for (int w = 0; w < 3; ++w) { + EXPECT_EQ(static_cast(macro_int_4_3_2_3[x][y][z][w]), 0); + } + } + } + } + auto macro_uint = FLAMEGPU->environment.getMacroProperty("macro_uint"); + EXPECT_EQ(static_cast(macro_uint), 0u); +} FLAMEGPU_STEP_FUNCTION(RESET_MACRO_ENV) { + validate_has_run = true; auto macro_float_3_3_3 = FLAMEGPU->environment.getMacroProperty("macro_float_3_3_3"); macro_float_3_3_3.zero(); auto macro_int_4_3_2_3 = FLAMEGPU->environment.getMacroProperty("macro_int_4_3_2_3"); @@ -140,6 +168,29 @@ FLAMEGPU_STEP_FUNCTION(RESET_MACRO_ENV) { auto macro_uint = FLAMEGPU->environment.getMacroProperty("macro_uint"); macro_uint.zero(); } +FLAMEGPU_STEP_FUNCTION(EXPORT_MACRO_ENV) { + FLAMEGPU->environment.exportMacroProperty("macro_float_3_3_3", current_test_file_name); +} +FLAMEGPU_STEP_FUNCTION(IMPORT_MACRO_ENV) { + const auto macro_float_3_3_3 = FLAMEGPU->environment.getMacroProperty("macro_float_3_3_3"); + int i = 0; + for (int x = 0; x < 3; ++x) { + for (int y = 0; y < 3; ++y) { + for (int z = 0; z < 3; ++z) { + EXPECT_EQ(static_cast(macro_float_3_3_3[x][y][z]), 0); + } + } + } + FLAMEGPU->environment.importMacroProperty("macro_float_3_3_3", current_test_file_name); + i = 0; + for (int x = 0; x < 3; ++x) { + for (int y = 0; y < 3; ++y) { + for (int z = 0; z < 3; ++z) { + EXPECT_EQ(static_cast(macro_float_3_3_3[x][y][z]), static_cast(i++)); + } + } + } +} class MiniSim { std::string test_file; @@ -650,5 +701,74 @@ TEST(IOTest2, AgentID_FileInput_IDCollision) { // Cleanup ASSERT_EQ(::remove(JSON_FILE_NAME), 0); } + + +class MiniSim3 { + std::string test_file; + + public: + ~MiniSim3() { + // Cleanup + if (!test_file.empty()) + ::remove(test_file.c_str()); + } + void run(const char *test_file_name) { + current_test_file_name = test_file_name; + this->test_file = test_file_name; + // Model description + ModelDescription model("test_model"); + AgentDescription a = model.newAgent("a"); + { + EnvironmentDescription e = model.Environment(); + e.newMacroProperty("macro_float_3_3_3"); + e.newMacroProperty("macro_int_4_3_2_3"); + e.newMacroProperty("macro_uint"); + } + AgentVector pop(a, 1); + model.addInitFunction(INIT_MACRO_ENV); + model.newLayer().addHostFunction(EXPORT_MACRO_ENV); + model.newLayer().addHostFunction(RESET_MACRO_ENV); + model.newLayer().addHostFunction(IMPORT_MACRO_ENV); + model.newLayer().addHostFunction(VALIDATE_MACRO_ENV2); + { // Validate env_vars + // Load model + CUDASimulation am(model); + // Step once, this checks and clears env vars + validate_has_run = false; + am.initFunctions(); + am.step(); + ASSERT_TRUE(validate_has_run); + } + } +}; + +class IOTest3 : public testing::Test { + protected: + void SetUp() override { + ms = new MiniSim3(); + } + + void TearDown() override { + delete ms; + } + + MiniSim3 *ms = nullptr; +}; + +TEST_F(IOTest3, XML_EnvMacroPropertyWriteRead) { + // Avoid fail if previous run didn't cleanup properly + ::remove(XML_FILE_NAME); + ms->run(XML_FILE_NAME); +} +TEST_F(IOTest3, JSON_EnvMacroPropertyWriteRead) { + // Avoid fail if previous run didn't cleanup properly + ::remove(JSON_FILE_NAME); + ms->run(JSON_FILE_NAME); +} +TEST_F(IOTest3, BIN_EnvMacroPropertyWriteRead) { + // Avoid fail if previous run didn't cleanup properly + ::remove(BIN_FILE_NAME); + ms->run(BIN_FILE_NAME); +} } // namespace test_io } // namespace flamegpu