diff --git a/include/flamegpu/gpu/CUDAAgent.h b/include/flamegpu/gpu/CUDAAgent.h index cff60e83d..5b867d7ff 100644 --- a/include/flamegpu/gpu/CUDAAgent.h +++ b/include/flamegpu/gpu/CUDAAgent.h @@ -208,10 +208,11 @@ class CUDAAgent : public AgentInterface { * * Uses Jitify to create an instantiation of the program. Any compilation errors in the user provided agent function will be reported here. * @param func The Agent function data structu containing the src for the function + * @param macro_env Object containing environment macro properties for the simulation instance * @param function_condition If true then this function will instantiate a function condition rather than an agent function * @throw exception::InvalidAgentFunc thrown if the user supplied agent function has compilation errors */ - void addInstantitateRTCFunction(const AgentFunctionData& func, bool function_condition = false); + void addInstantitateRTCFunction(const AgentFunctionData& func, const CUDAMacroEnvironment& macro_env, bool function_condition = false); /** * Returns the jitify kernel instantiation of the agent function. * Will throw an exception::InvalidAgentFunc excpetion if the function name does not have a valid instantiation diff --git a/include/flamegpu/gpu/CUDAMacroEnvironment.h b/include/flamegpu/gpu/CUDAMacroEnvironment.h new file mode 100644 index 000000000..b115d96e9 --- /dev/null +++ b/include/flamegpu/gpu/CUDAMacroEnvironment.h @@ -0,0 +1,262 @@ +#ifndef INCLUDE_FLAMEGPU_GPU_CUDAMACROENVIRONMENT_H_ +#define INCLUDE_FLAMEGPU_GPU_CUDAMACROENVIRONMENT_H_ + +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "detail/CUDAErrorChecking.cuh" +#include "flamegpu/runtime/detail/curve/curve.cuh" +#include "flamegpu/runtime/utility/HostMacroProperty.cuh" + +// forward declare classes from other modules + +namespace flamegpu { +namespace detail { +namespace curve { +class CurveRTCHost; +} // namespace curve +} // namespace detail + +struct SubEnvironmentData; +struct AgentFunctionData; +class EnvironmentDescription; +class CUDASimulation; + +/** + * This class is CUDASimulation's internal handler for macro environment functionality + */ +class CUDAMacroEnvironment { + /** + * This is the string used to generate MACRO_NAMESPACE_HASH + */ + static const char MACRO_NAMESPACE_STRING[18]; + /** + * Hash never changes, so we store a copy at creation + * Also ensure the device constexpr version matches + */ + const detail::curve::Curve::NamespaceHash MACRO_NAMESPACE_HASH; + /** + * Used to group items required by properties + */ + struct MacroEnvProp { + /** + * @param _type The type index of the base type (e.g. typeid(float)) + * @param _type_size The size of the base type (e.g. sizeof(float)) + * @param _elements Number of elements in each dimension + */ + MacroEnvProp(const std::type_index& _type, const size_t &_type_size, const std::array &_elements) + : type(_type) + , type_size(_type_size) + , elements(_elements) + , d_ptr(nullptr) + , is_sub(false) { } + ~MacroEnvProp() { + if (d_ptr && !is_sub) { + gpuErrchk(cudaFree(d_ptr)); + } + } + MacroEnvProp(const MacroEnvProp& other) = delete; + MacroEnvProp(MacroEnvProp&& other) + : type(other.type) + , type_size(other.type_size) + , elements(other.elements) + , d_ptr(other.d_ptr) + , is_sub(other.is_sub) { + other.d_ptr = nullptr; + } + std::type_index type; + size_t type_size; + std::array elements; + void *d_ptr; + // Denotes whether d_ptr is owned by this struct or not + bool is_sub; + // ptrdiff_t rtc_offset; // This is set by buildRTCOffsets(); + }; + const CUDASimulation& cudaSimulation; + std::map properties; + std::map> host_cache; + + public: + /** + * Normal constructor + * @param description Agent description of the agent + * @param _cudaSimulation Parent CUDASimulation of the agent + */ + CUDAMacroEnvironment(const EnvironmentDescription& description, const CUDASimulation& _cudaSimulation); + CUDAMacroEnvironment(CUDAMacroEnvironment&) = delete; + CUDAMacroEnvironment(CUDAMacroEnvironment&&) = delete; + /** + * Performs CUDA allocations, and registers CURVE variables + */ + void init(); + /** + * Performs CUDA allocations, and registers CURVE variables + * Initialises submodel mappings too + * @param mapping The SubEnvironment mapping info + * @param master_macro_env The master model's macro env to map sub macro properties with + * @note This must be called after the master model CUDAMacroEnvironment has init + */ + void init(const SubEnvironmentData& mapping, const CUDAMacroEnvironment& master_macro_env); + /** + * Release all CUDA allocations, and unregisters CURVE variables + */ + void free(); + /** + * Clears all CUDA pointers without deallocating, (e.g. if device has been reset) + */ + void purge(); + /** + * Register the properties to CURVE for use within the passed agent function + */ + void mapRuntimeVariables() const; + /** + * Release the properties from CURVE as registered for use within the passed agent function + */ + void unmapRuntimeVariables() const; + /** + * Register the properties to the provided RTC header + * @param curve_header The RTC header to act upon + */ + void mapRTCVariables(detail::curve::CurveRTCHost& curve_header) const; + /** + * Release the properties to the provided RTC header + * @param curve_header The RTC header to act upon + */ + void unmapRTCVariables(detail::curve::CurveRTCHost& curve_header) const; + +#if !defined(SEATBELTS) || SEATBELTS + /** + * Reset the flags used by seatbelts to catch potential race conditions + * @param streams Streams to async reset over + */ + void resetFlagsAsync(const std::vector& streams); + /** + * Returns the current state of the device read flag for the named macro property + * @param property_name Name of the macro property to query + */ + bool getDeviceReadFlag(const std::string& property_name); + /** + * Returns the current state of the device read flag for the named macro property + * @param property_name Name of the macro property to query + */ + bool getDeviceWriteFlag(const std::string& property_name); + /** + * Returns the raw flag variable used for read write flag for the named macro property + * @param property_name Name of the macro property to query + * @note This means both can be checked with a single memcpy + * @note Bit 0 being set means read has occurred, bit 1 being set means write has occurred + */ + unsigned int getDeviceRWFlags(const std::string& property_name); +#endif + /** + * Returns a HostAPI style direct accessor the macro property + * @param name Name of the macro property + * @tparam T Type of the macro property + * @tparam I 1st dimension length + * @tparam J 2nd dimension length + * @tparam K 3rd dimension length + * @tparam W 4th dimension length + */ + template + HostMacroProperty getProperty(const std::string& name); +#ifdef SWIG + /** + * Returns a HostAPI style direct accessor the macro property for SWIG + * @param name Name of the macro property + * @tparam T Type of the macro property + */ + template + HostMacroProperty_swig getProperty_swig(const std::string& name); +#endif +}; + +template +HostMacroProperty CUDAMacroEnvironment::getProperty(const std::string& name) { + // Validation + auto prop = properties.find(name); + if (prop == properties.end()) { + THROW flamegpu::exception::InvalidEnvProperty("Environment macro property with name '%s' not found, " + "in HostEnvironment::getMacroProperty()\n", + name.c_str()); + } else if (prop->second.type != std::type_index(typeid(T))) { + THROW flamegpu::exception::InvalidEnvProperty("Environment macro property '%s' type mismatch '%s' != '%s', " + "in HostEnvironment::getMacroProperty()\n", + name.c_str(), std::type_index(typeid(T)).name(), prop->second.type.name()); + } else if (prop->second.elements != std::array{I, J, K, W}) { + THROW flamegpu::exception::InvalidEnvProperty("Environment macro property '%s' dimensions mismatch (%u, %u, %u, %u) != (%u, %u, %u, %u), " + "in HostEnvironment::getMacroProperty()\n", + name.c_str(), I, J, K, W, prop->second.elements[0], prop->second.elements[1], prop->second.elements[2], prop->second.elements[3]); + } +#if !defined(SEATBELTS) || SEATBELTS + const unsigned int flags = getDeviceWriteFlag(name); + if (flags & (1 << 1)) { + THROW flamegpu::exception::InvalidOperation("Environment macro property '%s' was written to by an agent function in the same layer, " + "accessing it with a host function in the same layer could cause a race condition, in CUDAMacroEnvironment::getProperty().", + name.c_str()); + } + const bool read_flag = flags & (1 << 0); +#else + const bool read_flag = false; +#endif + // See if there is a live metadata in cache + auto cache = host_cache.find(name); + if (cache != host_cache.end()) { + if (cache->second.lock()) { + return HostMacroProperty(cache->second.lock()); + } + host_cache.erase(cache); + } + auto ret = std::make_shared(prop->second.d_ptr, prop->second.elements, sizeof(T), read_flag, name); + host_cache.emplace(name, ret); + return HostMacroProperty(ret); +} + +#ifdef SWIG +template +HostMacroProperty_swig CUDAMacroEnvironment::getProperty_swig(const std::string& name) { + // Validation + auto prop = properties.find(name); + if (prop == properties.end()) { + THROW flamegpu::exception::InvalidEnvProperty("Environment macro property with name '%s' not found, " + "in HostEnvironment::getMacroProperty()\n", + name.c_str()); + } else if (prop->second.type != std::type_index(typeid(T))) { + THROW flamegpu::exception::InvalidEnvProperty("Environment macro property '%s' type mismatch '%s' != '%s', " + "in HostEnvironment::getMacroProperty()\n", + name.c_str(), std::type_index(typeid(T)).name(), prop->second.type.name()); + } +#if !defined(SEATBELTS) || SEATBELTS + const unsigned int flags = getDeviceWriteFlag(name); + if (flags & (1 << 1)) { + THROW flamegpu::exception::InvalidOperation("Environment macro property '%s' was written to by an agent function in the same layer, " + "accessing it with a host function in the same layer could cause a race condition, in CUDAMacroEnvironment::getProperty().", + name.c_str()); + } + const bool read_flag = flags & (1 << 0); +#else + const bool read_flag = false; +#endif + // See if there is a live metadata in cache + auto cache = host_cache.find(name); + if (cache != host_cache.end()) { + if (cache->second.lock()) { + return HostMacroProperty_swig(cache->second.lock()); + } + host_cache.erase(cache); + } + auto ret = std::make_shared(prop->second.d_ptr, prop->second.elements, sizeof(T), read_flag, name); + host_cache.emplace(name, ret); + return HostMacroProperty_swig(ret); +} +#endif + +} // namespace flamegpu + +#endif // INCLUDE_FLAMEGPU_GPU_CUDAMACROENVIRONMENT_H_ diff --git a/include/flamegpu/gpu/CUDASimulation.h b/include/flamegpu/gpu/CUDASimulation.h index 4e6270e72..7e4529e1f 100644 --- a/include/flamegpu/gpu/CUDASimulation.h +++ b/include/flamegpu/gpu/CUDASimulation.h @@ -14,6 +14,7 @@ #include "flamegpu/gpu/CUDAEnsemble.h" #include "flamegpu/runtime/utility/RandomManager.cuh" #include "flamegpu/runtime/HostNewAgentAPI.h" +#include "flamegpu/gpu/CUDAMacroEnvironment.h" #ifdef VISUALISATION #include "flamegpu/visualiser/ModelVis.h" @@ -332,6 +333,10 @@ class CUDASimulation : public Simulation { * Map of agent storage */ CUDAAgentMap agent_map; + /** + * Macro env property storage + */ + CUDAMacroEnvironment macro_env; /** * Internal model config */ diff --git a/include/flamegpu/model/EnvironmentDescription.h b/include/flamegpu/model/EnvironmentDescription.h index 96c6743f4..02915978e 100644 --- a/include/flamegpu/model/EnvironmentDescription.h +++ b/include/flamegpu/model/EnvironmentDescription.h @@ -58,17 +58,60 @@ class EnvironmentDescription { bool isConst; const util::Any data; bool operator==(const PropData &rhs) const { + if (this == &rhs) + return true; if (this->isConst != rhs.isConst || this->data.elements != rhs.data.elements || this->data.length != rhs.data.length || this->data.type != rhs.data.type) return false; + if (this->data.ptr == rhs.data.ptr) + return true; for (size_t i = 0; i < this->data.length; ++i) { if (static_cast(this->data.ptr)[i] != static_cast(rhs.data.ptr)[i]) return false; } return true; } + bool operator!=(const PropData& rhs) const { + return !operator==(rhs); + } + }; + /** + * Holds all of the properties required to add a value to EnvironmentManager + */ + struct MacroPropData { + /** + * @param _type The type index of the base type (e.g. typeid(float)) + * @param _type_size The size of the base type (e.g. sizeof(float)) + * @param _elements Number of elements in each dimension + */ + MacroPropData(const std::type_index &_type, const size_t _type_size, const std::array &_elements) + : type(_type) + , type_size(_type_size) + , elements(_elements) { } + std::type_index type; + size_t type_size; + std::array elements; + bool operator==(const MacroPropData& rhs) const { + if (this == &rhs) + return true; + if (this->type != rhs.type + || this->type_size != rhs.type_size + || this->elements[0] != rhs.elements[0] + || this->elements[1] != rhs.elements[1] + || this->elements[2] != rhs.elements[2] + || this->elements[3] != rhs.elements[3]) + return false; + for (size_t i = 0; i < this->elements.size(); ++i) { + if (this->elements[i] != rhs.elements[i]) + return false; + } + return true; + } + bool operator!=(const MacroPropData& rhs) const { + return !operator==(rhs); + } }; /** * Default destruction @@ -110,6 +153,38 @@ class EnvironmentDescription { */ template void newPropertyArray(const std::string &name, const EnvironmentManager::size_type &length, const std::vector &value, const bool &isConst = false); +#endif + /** + * Define a new environment macro property + * + * Environment macro properties are designed for large environment properties, too large of fast constant memory. + * This means they must instead be stored in slower global memory, however that allows them to be modified during agent functions via a limited set of atomic operations. + * + * @param name Name of the macro property + * @tparam T Type of the macro property + * @tparam I Length of the first dimension of the macro property, default 1 + * @tparam J Length of the second dimension of the macro property, default 1 + * @tparam K Length of the third dimension of the macro property, default 1 + * @tparam W Length of the fourth dimension of the macro property, default 1 + */ + template + void newMacroProperty(const std::string& name); +#ifdef SWIG + /** + * Define a new environment macro property, swig specific version + * + * Environment macro properties are designed for large environment properties, too large of fast constant memory. + * This means they must instead be stored in slower global memory, however that allows them to be modified during agent functions via a limited set of atomic operations. + * + * @param name Name of the macro property + * @param I Length of the first dimension of the macro property, default 1 + * @param J Length of the second dimension of the macro property, default 1 + * @param K Length of the third dimension of the macro property, default 1 + * @param W Length of the fourth dimension of the macro property, default 1 + * @tparam T Type of the macro property + */ + template + void newMacroProperty_swig(const std::string& name, EnvironmentManager::size_type I = 1, EnvironmentManager::size_type J = 1, EnvironmentManager::size_type K = 1, EnvironmentManager::size_type W = 1); #endif /** * Gets an environment property @@ -203,6 +278,7 @@ class EnvironmentDescription { #endif const std::unordered_map getPropertiesMap() const; + const std::unordered_map getMacroPropertiesMap() const; private: /** @@ -219,6 +295,10 @@ class EnvironmentDescription { * Main storage of all properties */ std::unordered_map properties{}; + /** + * Main storage of all macroproperties + */ + std::unordered_map macro_properties{}; }; @@ -229,7 +309,7 @@ template void EnvironmentDescription::newProperty(const std::string &name, const T &value, const bool &isConst) { if (!name.empty() && name[0] == '_') { THROW exception::ReservedName("Environment property names cannot begin with '_', this is reserved for internal usage, " - "in EnvironmentDescription::add()."); + "in EnvironmentDescription::newProperty()."); } // Limited to Arithmetic types // Compound types would allow host pointers inside structs to be passed @@ -237,7 +317,7 @@ void EnvironmentDescription::newProperty(const std::string &name, const T &value "Only arithmetic types can be used as environmental properties"); if (properties.find(name) != properties.end()) { THROW exception::DuplicateEnvProperty("Environmental property with name '%s' already exists, " - "in EnvironmentDescription::add().", + "in EnvironmentDescription::newProperty().", name.c_str()); } newProperty(name, reinterpret_cast(&value), sizeof(T), isConst, 1, typeid(T)); @@ -246,7 +326,7 @@ template void EnvironmentDescription::newProperty(const std::string &name, const std::array &value, const bool &isConst) { if (!name.empty() && name[0] == '_') { THROW exception::ReservedName("Environment property names cannot begin with '_', this is reserved for internal usage, " - "in EnvironmentDescription::add()."); + "in EnvironmentDescription::newProperty()."); } static_assert(N > 0, "Environment property arrays must have a length greater than 0."); // Limited to Arithmetic types @@ -255,7 +335,7 @@ void EnvironmentDescription::newProperty(const std::string &name, const std::arr "Only arithmetic types can be used as environmental properties"); if (properties.find(name) != properties.end()) { THROW exception::DuplicateEnvProperty("Environmental property with name '%s' already exists, " - "in EnvironmentDescription::add().", + "in EnvironmentDescription::newProperty().", name.c_str()); } newProperty(name, reinterpret_cast(value.data()), N * sizeof(T), isConst, N, typeid(T)); @@ -265,15 +345,15 @@ template void EnvironmentDescription::newPropertyArray(const std::string &name, const EnvironmentManager::size_type &N, const std::vector &value, const bool& isConst) { if (!name.empty() && name[0] == '_') { THROW exception::ReservedName("Environment property names cannot begin with '_', this is reserved for internal usage, " - "in EnvironmentDescription::addArray()."); + "in EnvironmentDescription::newPropertyArray()."); } if (value.size() != N) { THROW exception::InvalidEnvProperty("Environment property array length does not match the value provided, %u != %llu," - "in EnvironmentDescription::addArray().", N, value.size()); + "in EnvironmentDescription::newPropertyArray().", N, value.size()); } if (N == 0) { THROW exception::InvalidEnvProperty("Environment property arrays must have a length greater than 0." - "in EnvironmentDescription::addArray()."); + "in EnvironmentDescription::newPropertyArray()."); } // Limited to Arithmetic types // Compound types would allow host pointers inside structs to be passed @@ -281,7 +361,7 @@ void EnvironmentDescription::newPropertyArray(const std::string &name, const Env "Only arithmetic types can be used as environmental properties"); if (properties.find(name) != properties.end()) { THROW exception::DuplicateEnvProperty("Environmental property with name '%s' already exists, " - "in EnvironmentDescription::addArray().", + "in EnvironmentDescription::newPropertyArray().", name.c_str()); } newProperty(name, reinterpret_cast(value.data()), N * sizeof(T), isConst, N, typeid(T)); @@ -300,13 +380,13 @@ T EnvironmentDescription::getProperty(const std::string &name) const { if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } return *reinterpret_cast(i->second.data.ptr); } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", name.c_str()); } template @@ -319,12 +399,12 @@ std::array EnvironmentDescription::getProperty(const std::string &name) co if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property array ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } if (i->second.data.elements != N) { THROW exception::OutOfBoundsException("Length of named environmental property array (%u) does not match template argument N (%u), " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", i->second.data.elements, N); } // Copy old data to return @@ -333,7 +413,7 @@ std::array EnvironmentDescription::getProperty(const std::string &name) co return rtn; } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", name.c_str()); } template @@ -346,19 +426,19 @@ T EnvironmentDescription::getProperty(const std::string &name, const Environment if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property array ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } if (i->second.data.elements <= index) { THROW exception::OutOfBoundsException("Index (%u) exceeds named environmental property array's length (%u), " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", index, i->second.data.elements); } // Copy old data to return return *(reinterpret_cast(i->second.data.ptr) + index); } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::get().", + "in EnvironmentDescription::getProperty().", name.c_str()); } #ifdef SWIG @@ -372,7 +452,7 @@ std::vector EnvironmentDescription::getPropertyArray(const std::string& name) if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property array ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::getArray().", + "in EnvironmentDescription::getPropertyArray().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } // Copy old data to return @@ -381,7 +461,7 @@ std::vector EnvironmentDescription::getPropertyArray(const std::string& name) return rtn; } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::getArray().", + "in EnvironmentDescription::getPropertyArray().", name.c_str()); } #endif @@ -393,7 +473,7 @@ template T EnvironmentDescription::setProperty(const std::string &name, const T &value) { if (!name.empty() && name[0] == '_') { THROW exception::ReservedName("Environment property names cannot begin with '_', this is reserved for internal usage, " - "in EnvironmentDescription::set()."); + "in EnvironmentDescription::setProperty()."); } // Limited to Arithmetic types // Compound types would allow host pointers inside structs to be passed @@ -403,7 +483,7 @@ T EnvironmentDescription::setProperty(const std::string &name, const T &value) { if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } // Copy old data to return @@ -413,14 +493,14 @@ T EnvironmentDescription::setProperty(const std::string &name, const T &value) { return rtn; } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", name.c_str()); } template std::array EnvironmentDescription::setProperty(const std::string &name, const std::array &value) { if (!name.empty() && name[0] == '_') { THROW exception::ReservedName("Environment property names cannot begin with '_', this is reserved for internal usage, " - "in EnvironmentDescription::set()."); + "in EnvironmentDescription::setProperty()."); } // Limited to Arithmetic types // Compound types would allow host pointers inside structs to be passed @@ -430,12 +510,12 @@ std::array EnvironmentDescription::setProperty(const std::string &name, co if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property array ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } if (i->second.data.elements != N) { THROW exception::OutOfBoundsException("Length of named environmental property array (%u) does not match template argument N (%u), " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", i->second.data.elements, N); } // Copy old data to return @@ -446,14 +526,14 @@ std::array EnvironmentDescription::setProperty(const std::string &name, co return rtn; } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", name.c_str()); } template T EnvironmentDescription::setProperty(const std::string &name, const EnvironmentManager::size_type &index, const T &value) { if (!name.empty() && name[0] == '_') { THROW exception::ReservedName("Environment property names cannot begin with '_', this is reserved for internal usage, " - "in EnvironmentDescription::set()."); + "in EnvironmentDescription::setProperty()."); } // Limited to Arithmetic types // Compound types would allow host pointers inside structs to be passed @@ -463,12 +543,12 @@ T EnvironmentDescription::setProperty(const std::string &name, const Environment if (i != properties.end()) { if (i->second.data.type != std::type_index(typeid(T))) { THROW exception::InvalidEnvPropertyType("Environmental property array ('%s') type (%s) does not match template argument T (%s), " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", name.c_str(), i->second.data.type.name(), typeid(T).name()); } if (i->second.data.elements <= index) { THROW exception::OutOfBoundsException("Index (%u) exceeds named environmental property array's length (%u), " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", index, i->second.data.elements); } // Copy old data to return @@ -478,7 +558,7 @@ T EnvironmentDescription::setProperty(const std::string &name, const Environment return rtn; } THROW exception::InvalidEnvProperty("Environmental property with name '%s' does not exist, " - "in EnvironmentDescription::set().", + "in EnvironmentDescription::setProperty().", name.c_str()); } #ifdef SWIG @@ -516,6 +596,62 @@ std::vector EnvironmentDescription::setPropertyArray(const std::string& name, name.c_str()); } #endif +template +void EnvironmentDescription::newMacroProperty(const std::string& name) { + if (!name.empty() && name[0] == '_') { + THROW exception::ReservedName("Environment macro property names cannot begin with '_', this is reserved for internal usage, " + "in EnvironmentDescription::newMacroProperty()."); + } + // Limited to Arithmetic types + // Compound types would allow host pointers inside structs to be passed + static_assert(std::is_arithmetic::value || std::is_enum::value, + "Only arithmetic types can be used as environmental macro properties"); + static_assert(I > 0, "Environment macro properties must have a length greater than 0 in the first axis."); + static_assert(J > 0, "Environment macro properties must have a length greater than 0 in the second axis."); + static_assert(K > 0, "Environment macro properties must have a length greater than 0 in the third axis."); + static_assert(W > 0, "Environment macro properties must have a length greater than 0 in the fourth axis."); + if (macro_properties.find(name) != macro_properties.end()) { + THROW exception::DuplicateEnvProperty("Environmental macro property with name '%s' already exists, " + "in EnvironmentDescription::newMacroProperty().", + name.c_str()); + } + macro_properties.emplace(name, MacroPropData(typeid(T), sizeof(T), { I, J, K, W })); +} +#ifdef SWIG +template +void EnvironmentDescription::newMacroProperty_swig(const std::string& name, EnvironmentManager::size_type I, EnvironmentManager::size_type J, EnvironmentManager::size_type K, EnvironmentManager::size_type W) { + if (!name.empty() && name[0] == '_') { + THROW exception::ReservedName("Environment macro property names cannot begin with '_', this is reserved for internal usage, " + "in EnvironmentDescription::newMacroProperty()."); + } + // Limited to Arithmetic types + // Compound types would allow host pointers inside structs to be passed + static_assert(std::is_arithmetic::value || std::is_enum::value, + "Only arithmetic types can be used as environmental macro properties"); + if (I <= 0) { + THROW exception::DuplicateEnvProperty("Environmental macro property with name '%s' must have a length greater than 0 in the first axis, " + "in EnvironmentDescription::newMacroProperty().", + name.c_str()); + } else if (J <= 0) { + THROW exception::DuplicateEnvProperty("Environmental macro property with name '%s' must have a length greater than 0 in the second axis, " + "in EnvironmentDescription::newMacroProperty().", + name.c_str()); + } else if (K <= 0) { + THROW exception::DuplicateEnvProperty("Environmental macro property with name '%s' must have a length greater than 0 in the third axis, " + "in EnvironmentDescription::newMacroProperty().", + name.c_str()); + } else if (W <= 0) { + THROW exception::DuplicateEnvProperty("Environmental macro property with name '%s' must have a length greater than 0 in the fourth axis, " + "in EnvironmentDescription::newMacroProperty().", + name.c_str()); + } else if (macro_properties.find(name) != macro_properties.end()) { + THROW exception::DuplicateEnvProperty("Environmental macro property with name '%s' already exists, " + "in EnvironmentDescription::newMacroProperty().", + name.c_str()); + } + macro_properties.emplace(name, MacroPropData(typeid(T), sizeof(T), { I, J, K, W })); +} +#endif } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_MODEL_ENVIRONMENTDESCRIPTION_H_ diff --git a/include/flamegpu/model/SubEnvironmentData.h b/include/flamegpu/model/SubEnvironmentData.h index e1fd33f56..0afd9214a 100644 --- a/include/flamegpu/model/SubEnvironmentData.h +++ b/include/flamegpu/model/SubEnvironmentData.h @@ -46,6 +46,10 @@ struct SubEnvironmentData : std::enable_shared_from_this { * Holds all of the model's environment property mappings */ Mapping properties; + /** + * Holds all of the model's environment macro property mappings + */ + Mapping macro_properties; /** * The model which this environment is a member of */ diff --git a/include/flamegpu/model/SubEnvironmentDescription.h b/include/flamegpu/model/SubEnvironmentDescription.h index 0ace17eb1..c9e8e25ad 100644 --- a/include/flamegpu/model/SubEnvironmentDescription.h +++ b/include/flamegpu/model/SubEnvironmentDescription.h @@ -42,12 +42,23 @@ class SubEnvironmentDescription { SubEnvironmentDescription& operator=(SubEnvironmentDescription &&other_agent) noexcept = delete; public: + /** + * Automatically map all compatible properties and macro properties + * In order to be compatible, properties must share the same name, type, dimensions/length (number of elements) + * Const master properties cannot be mapped to non-const sub properties, however the inverse is permitted + */ + void autoMap(); /** * Automatically map all compatible properties * In order to be compatible, properties must share the same name, type, length (number of elements) * Const master properties cannot be mapped to non-const sub properties, however the inverse is permitted */ void autoMapProperties(); + /** + * Automatically map all compatible macro properties + * In order to be compatible, properties must share the same name, type, dimensions + */ + void autoMapMacroProperties(); /** * Links the named properties between the master and sub environment * In order to be compatible, properties must share the same name, type, length (number of elements) @@ -60,12 +71,29 @@ class SubEnvironmentDescription { */ void mapProperty(const std::string &sub_property_name, const std::string &master_property_name); /** - * Returns the master agent property which has been mapped to the name subagent state + * Links the named macro properties between the master and sub environment + * In order to be compatible, macro properties must share the same name, type, dimensions + * @param sub_property_name Name of the macro property in the sub models agent + * @param master_property_name Name of the macro property in the master models agent + * @throws exception::InvalidParent If the sub agent or master agent weak_ptrs have expired (this should never happen) + * @throws exception::InvalidEnvProperty If the named macro property does not exist within the bound sub or master environment + * @throws exception::InvalidEnvProperty If the named macro properties do not share the same type and length + */ + void mapMacroProperty(const std::string& sub_property_name, const std::string& master_property_name); + /** + * Returns the name of the master property which has been mapped to the named subenvironment property * @param sub_property_name Name of the state in the sub agent to check * @return The name of the state within the master agent which is mapped * @throws exception::InvalidEnvProperty If the sub environment property does not exist or has not been mapped yet */ std::string getPropertyMapping(const std::string &sub_property_name); + /** + * Returns the name of the master macro property which has been mapped to the named subenvironment macro property + * @param sub_property_name Name of the state in the sub agent to check + * @return The name of the state within the master agent which is mapped + * @throws exception::InvalidEnvProperty If the sub environment property does not exist or has not been mapped yet + */ + std::string getMacroPropertyMapping(const std::string& sub_property_name); private: /** diff --git a/include/flamegpu/runtime/DeviceAPI.cuh b/include/flamegpu/runtime/DeviceAPI.cuh index a89a601bf..4fd62a766 100644 --- a/include/flamegpu/runtime/DeviceAPI.cuh +++ b/include/flamegpu/runtime/DeviceAPI.cuh @@ -88,6 +88,14 @@ class ReadOnlyDeviceAPI { return getVariable("_id"); } + /** + * Access the current stepCount + * @return the current step count, 0 indexed unsigned. + */ + __forceinline__ __device__ unsigned int getStepCounter() const { + return environment.getProperty("_stepCount"); + } + /** * Provides access to random functionality inside agent functions * @note random state isn't stored within the object, so it can be const @@ -96,15 +104,7 @@ class ReadOnlyDeviceAPI { /** * Provides access to environment variables inside agent functions */ - const DeviceEnvironment environment; - - /** - * Access the current stepCount - * @return the current step count, 0 indexed unsigned. - */ - __forceinline__ __device__ unsigned int getStepCounter() const { - return environment.getProperty("_stepCount"); - } + const ReadOnlyDeviceEnvironment environment; /** * Returns the current CUDA thread of the agent @@ -144,7 +144,7 @@ class ReadOnlyDeviceAPI { * @tparam MessageOut Output message type (the form found in flamegpu/runtime/messaging.h, MessageNone etc) */ template -class DeviceAPI : public ReadOnlyDeviceAPI{ +class DeviceAPI { // Friends have access to TID() & TS_ID() template friend __global__ void agent_function_wrapper( @@ -259,11 +259,36 @@ class DeviceAPI : public ReadOnlyDeviceAPI{ unsigned int *&scanFlag_agentOutput, typename MessageIn::In &&message_in, typename MessageOut::Out &&message_out) - : ReadOnlyDeviceAPI(instance_id_hash, agentfuncname_hash, d_rng) - , message_in(message_in) + : message_in(message_in) , message_out(message_out) , agent_out(AgentOut(_agent_output_hash, d_agent_output_nextID, scanFlag_agentOutput)) + , random(AgentRandom(&d_rng[getThreadIndex()])) + , environment(DeviceEnvironment(instance_id_hash)) + , agent_func_name_hash(agentfuncname_hash) { } + /** + * Returns the specified variable from the currently executing agent + * @param variable_name name used for accessing the variable, this value should be a string literal e.g. "foobar" + * @tparam T Type of the agent variable being accessed + * @tparam N Length of variable name, this should always be implicit if passing a string literal + * @throws exception::DeviceError If name is not a valid variable within the agent (flamegpu must be built with SEATBELTS enabled for device error checking) + * @throws exception::DeviceError If T is not the type of variable 'name' within the agent (flamegpu must be built with SEATBELTS enabled for device error checking) + */ + template __device__ + T getVariable(const char(&variable_name)[N]) const; + /** + * Returns the specified variable array element from the currently executing agent + * @param variable_name name used for accessing the variable, this value should be a string literal e.g. "foobar" + * @param index Index of the element within the variable array to return + * @tparam T Type of the agent variable being accessed + * @tparam N The length of the array variable, as set within the model description hierarchy + * @tparam M Length of variable_name, this should always be implicit if passing a string literal + * @throws exception::DeviceError If name is not a valid variable within the agent (flamegpu must be built with SEATBELTS enabled for device error checking) + * @throws exception::DeviceError If T is not the type of variable 'name' within the agent (flamegpu must be built with SEATBELTS enabled for device error checking) + * @throws exception::DeviceError If index is out of bounds for the variable array specified by name (flamegpu must be built with SEATBELTS enabled for device error checking) + */ + template __device__ + T getVariable(const char(&variable_name)[M], const unsigned int &index) const; /** * Sets a variable within the currently executing agent * @param variable_name The name of the variable @@ -289,6 +314,44 @@ class DeviceAPI : public ReadOnlyDeviceAPI{ */ template __device__ void setVariable(const char(&variable_name)[M], const unsigned int &index, const T &value); + /** + * Returns the agent's unique identifier + */ + __device__ id_t getID() { + return getVariable("_id"); + } + + /** + * Access the current stepCount + * @return the current step count, 0 indexed unsigned. + */ + __forceinline__ __device__ unsigned int getStepCounter() const { + return environment.getProperty("_stepCount"); + } + + /** + * Returns the current CUDA thread of the agent + * All agents execute in a unique thread, but their associated thread may change between agent functions + * Thread indices begin at 0 and continue to 1 below the number of agents executing + */ + __forceinline__ __device__ static unsigned int getThreadIndex() { + /* + // 3D version + auto blockId = blockIdx.x + blockIdx.y * gridDim.x + + gridDim.x * gridDim.y * blockIdx.z; + auto threadId = blockId * (blockDim.x * blockDim.y * blockDim.z) + + (threadIdx.z * (blockDim.x * blockDim.y)) + + (threadIdx.y * blockDim.x) + + threadIdx.x; + return threadId;*/ +#ifdef SEATBELTS + assert(blockDim.y == 1); + assert(blockDim.z == 1); + assert(gridDim.y == 1); + assert(gridDim.z == 1); +#endif + return blockIdx.x * blockDim.x + threadIdx.x; + } /** * Provides access to message read functionality inside agent functions @@ -302,6 +365,18 @@ class DeviceAPI : public ReadOnlyDeviceAPI{ * Provides access to agent output functionality inside agent functions */ const AgentOut agent_out; + /** + * Provides access to random functionality inside agent functions + * @note random state isn't stored within the object, so it can be const + */ + const AgentRandom random; + /** + * Provides access to environment variables inside agent functions + */ + const DeviceEnvironment environment; + + protected: + detail::curve::Curve::NamespaceHash agent_func_name_hash; }; @@ -318,21 +393,34 @@ __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[N]) const // return the variable from curve return value; } +template +__device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const unsigned int &array_index) const { + // simple indexing assumes index is the thread number (this may change later) + const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; + + // get the value from curve + T value = detail::curve::Curve::getAgentArrayVariable(variable_name, agent_func_name_hash , index, array_index); + + // return the variable from curve + return value; +} template template -__device__ void DeviceAPI::setVariable(const char(&variable_name)[N], T value) { - if (variable_name[0] == '_') { - return; // Fail silently - } +__device__ T DeviceAPI::getVariable(const char(&variable_name)[N]) const { // simple indexing assumes index is the thread number (this may change later) const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; - // set the variable using curve - detail::curve::Curve::setAgentVariable(variable_name, agent_func_name_hash, value, index); + + // get the value from curve + T value = detail::curve::Curve::getAgentVariable(variable_name, agent_func_name_hash , index); + + // return the variable from curve + return value; } +template template -__device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const unsigned int &array_index) const { +__device__ T DeviceAPI::getVariable(const char(&variable_name)[M], const unsigned int &array_index) const { // simple indexing assumes index is the thread number (this may change later) const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; @@ -343,6 +431,17 @@ __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const return value; } +template +template +__device__ void DeviceAPI::setVariable(const char(&variable_name)[N], T value) { + if (variable_name[0] == '_') { + return; // Fail silently + } + // simple indexing assumes index is the thread number (this may change later) + const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; + // set the variable using curve + detail::curve::Curve::setAgentVariable(variable_name, agent_func_name_hash, value, index); +} template template __device__ void DeviceAPI::setVariable(const char(&variable_name)[M], const unsigned int &array_index, const T &value) { diff --git a/include/flamegpu/runtime/HostAPI.h b/include/flamegpu/runtime/HostAPI.h index 7e09c5828..e4608b16d 100644 --- a/include/flamegpu/runtime/HostAPI.h +++ b/include/flamegpu/runtime/HostAPI.h @@ -19,6 +19,7 @@ namespace flamegpu { class CUDAScatter; class CUDASimulation; class HostAgentAPI; +class CUDAMacroEnvironment; /** * @brief A flame gpu api class for use by host functions only @@ -51,6 +52,7 @@ class HostAPI { CUDAScatter &scatter, const AgentOffsetMap &agentOffsets, AgentDataMap &agentData, + CUDAMacroEnvironment ¯o_env, const unsigned int &streamId, cudaStream_t stream); /** diff --git a/include/flamegpu/runtime/detail/curve/curve_rtc.cuh b/include/flamegpu/runtime/detail/curve/curve_rtc.cuh index a5cbadc9c..c830744e1 100644 --- a/include/flamegpu/runtime/detail/curve/curve_rtc.cuh +++ b/include/flamegpu/runtime/detail/curve/curve_rtc.cuh @@ -1,6 +1,7 @@ #ifndef INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_RTC_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_RTC_CUH_ +#include #include #include #include @@ -142,6 +143,22 @@ class CurveRTCHost { * @throws exception::UnknownInternalError If the specified property is not registered */ void unregisterEnvVariable(const char* propertyName); + /** + * Specify an environment macro property to be included in the dynamic header + * @param propertyName The property's name + * @param d_ptr Pointer to the buffer in device memory + * @param type The name of the property's type (%std::type_index::name()) + * @param type_size The type size of the property's base type (sizeof()), this is the size of a single element if the property is an array property. + * @param dimensions The number of elements in the property (1 unless the property is an array property) + * @throws exception::UnknownInternalError If an environment property with the same name is already registered + */ + void registerEnvMacroProperty(const char* propertyName, void* d_ptr, const char* type, size_t type_size, const std::array& dimensions); + /** + * Unregister an environment property, so that it is nolonger included in the dynamic header + * @param propertyName The property's name + * @throws exception::UnknownInternalError If the specified property is not registered + */ + void unregisterEnvMacroProperty(const char* propertyName); /** * Set the filename tagged in the file (goes into a #line statement) * @param filename Name to be used for the file in compile errors @@ -243,6 +260,32 @@ class CurveRTCHost { */ size_t type_size; }; + /** + * Properties for a registered environment macro property + */ + struct RTCEnvMacroPropertyProperties { + /** + * Name of the property's base type (e.g. type of an individual element if array property) + */ + std::string type; + /** + * Number of elemements in each dimension + */ + std::array dimensions; + /** + * Size of the property's base type (e.g. size of an individual element if array property) + */ + size_t type_size; + /** + * Copy of the device pointer + * @note This assumes it will never be reallocated/resized after registration + */ + void* d_ptr; + /** + * Pointer to a location in host memory where the device pointer to this variables buffer must be stored + */ + void* h_data_ptr; + }; private: /** @@ -288,6 +331,10 @@ class CurveRTCHost { * Offset into h_data_buffer where output agent (device agent birth) variable data begins */ size_t newAgent_data_offset = 0; + /** + * Offset into h_data_buffer where output agent (device agent birth) variable data begins + */ + size_t envMacro_data_offset = 0; /** * Size of the allocation pointed to by h_data_buffer */ @@ -321,6 +368,11 @@ class CurveRTCHost { * */ std::map RTCEnvVariables; + /** + * Registered environment macro property properties + * + */ + std::map RTCEnvMacroProperties; }; } // namespace curve diff --git a/include/flamegpu/runtime/utility/DeviceEnvironment.cuh b/include/flamegpu/runtime/utility/DeviceEnvironment.cuh index b693ca722..64d1c99bd 100644 --- a/include/flamegpu/runtime/utility/DeviceEnvironment.cuh +++ b/include/flamegpu/runtime/utility/DeviceEnvironment.cuh @@ -2,10 +2,11 @@ #define INCLUDE_FLAMEGPU_RUNTIME_UTILITY_DEVICEENVIRONMENT_CUH_ // #include -#include #include #include +#include "flamegpu/runtime/utility/DeviceMacroProperty.cuh" + namespace flamegpu { #ifndef __CUDACC_RTC__ @@ -17,12 +18,14 @@ namespace detail { } // namespace detail #endif + + /** * Utility for accessing environmental properties * These can only be read within agent functions * They can be set and updated within host functions */ -class DeviceEnvironment { +class ReadOnlyDeviceEnvironment { /** * Constructs the object */ @@ -31,11 +34,19 @@ class DeviceEnvironment { * Performs runtime validation that CURVE_NAMESPACE_HASH matches host value */ friend class EnvironmentManager; + friend class CUDAMacroEnvironment; + + protected: /** * Device accessible copy of curve namespace hash, this is precomputed from EnvironmentManager::CURVE_NAMESPACE_HASH - * EnvironmentManager::EnvironmentManager() validates that this value matches + * EnvironmentManager::initialiseDevice() validates that this value matches */ __host__ __device__ static constexpr unsigned int CURVE_NAMESPACE_HASH() { return 0X1428F902u; } + /** + * Device accessible copy of macro namespace hash, this is precomputed from CUDAMacroEnvironment::MACRO_NAMESPACE_HASH + * CUDAMacroEnvironment::CUDAMacroEnvironment() validates that this value matches + */ + __host__ __device__ static constexpr unsigned int MACRO_NAMESPACE_HASH() { return 0xF3ABEB4F; } /** * Hash of the model's name, this is added to CURVE_NAMESPACE_HASH and variable name hash to find curve hash */ @@ -44,7 +55,7 @@ class DeviceEnvironment { * Constructor, requires the model name hash to init modelname_hash * @param _modelname_hash Hash of model name generated by curveGetVariableHash() */ - __device__ __forceinline__ DeviceEnvironment(const detail::curve::Curve::NamespaceHash &_modelname_hash) + __device__ __forceinline__ ReadOnlyDeviceEnvironment(const detail::curve::Curve::NamespaceHash &_modelname_hash) : modelname_hash(_modelname_hash) { } public: @@ -78,6 +89,44 @@ class DeviceEnvironment { */ template __device__ __forceinline__ bool containsProperty(const char(&name)[N]) const; + /** + * Returns a read-only accessor to the named macro property + * @param name name used for accessing the property, this value should be a string literal e.g. "foobar" + * @tparam I Length of macro property in the 1st dimension, default 1 + * @tparam J Length of macro property in the 2nd dimension, default 1 + * @tparam K Length of macro property in the 3rd dimension, default 1 + * @tparam W Length of macro property in the 4th dimension, default 1 + * @tparam N Length of variable name, this should always be implicit if passing a string literal + */ + template + __device__ __forceinline__ ReadOnlyDeviceMacroProperty getMacroProperty(const char(&name)[N]) const; +}; +/** + * Utility for accessing environmental properties + * These can only be read within agent functions + * They can be set and updated within host functions + * This version also allows limited write access to device macro properties + */ +class DeviceEnvironment : public ReadOnlyDeviceEnvironment { + public: + /** + * Constructor, requires the model name hash to init modelname_hash + * @param _modelname_hash Hash of model name generated by curveGetVariableHash() + * @note Public because it's constructed by a templated class + */ + __device__ __forceinline__ DeviceEnvironment(const detail::curve::Curve::NamespaceHash& _modelname_hash) + : ReadOnlyDeviceEnvironment(_modelname_hash) { } + /** + * Returns a read-only accessor to the named macro property + * @param name name used for accessing the property, this value should be a string literal e.g. "foobar" + * @tparam I Length of macro property in the 1st dimension, default 1 + * @tparam J Length of macro property in the 2nd dimension, default 1 + * @tparam K Length of macro property in the 3rd dimension, default 1 + * @tparam W Length of macro property in the 4th dimension, default 1 + * @tparam N Length of variable name, this should always be implicit if passing a string literal + */ + template + __device__ __forceinline__ DeviceMacroProperty getMacroProperty(const char(&name)[N]) const; }; // Mash compilation of these functions from RTC builds as this requires a dynamic implementation of the function in curve_rtc @@ -86,7 +135,7 @@ class DeviceEnvironment { * Getters */ template -__device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N]) const { +__device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[N]) const { detail::curve::Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); const auto cv = detail::curve::Curve::getVariable(cvh); #if !defined(SEATBELTS) || SEATBELTS @@ -108,7 +157,7 @@ __device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N] #endif } template -__device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N], const unsigned int &index) const { +__device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[N], const unsigned int &index) const { detail::curve::Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); const auto cv = detail::curve::Curve::getVariable(cvh); #if !defined(SEATBELTS) || SEATBELTS @@ -131,11 +180,51 @@ __device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N] * Util */ template -__device__ __forceinline__ bool DeviceEnvironment::containsProperty(const char(&name)[N]) const { +__device__ __forceinline__ bool ReadOnlyDeviceEnvironment::containsProperty(const char(&name)[N]) const { detail::curve::Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); return detail::curve::Curve::getVariable(cvh) != detail::curve::Curve::UNKNOWN_VARIABLE; } +template +__device__ __forceinline__ ReadOnlyDeviceMacroProperty ReadOnlyDeviceEnvironment::getMacroProperty(const char(&name)[N]) const { + detail::curve::Curve::VariableHash cvh = MACRO_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); + const auto cv = detail::curve::Curve::getVariable(cvh); +#if !defined(SEATBELTS) || SEATBELTS + if (cv == detail::curve::Curve::UNKNOWN_VARIABLE) { + DTHROW("Environment macro property name: %s was not found.\n", name); + } else if (detail::curve::detail::d_sizes[cv] != sizeof(T)) { + DTHROW("Environment macro property with name: %s type size mismatch %llu != %llu.\n", name, detail::curve::detail::d_sizes[cv], sizeof(T)); + } else if (detail::curve::detail::d_lengths[cv] != I * J * K * W) { + DTHROW("Environment macro property with name: %s total length mismatch (%u != %u).\n", name, detail::curve::detail::d_lengths[cv], I * J * K * W); + } else { + return DeviceMacroProperty(reinterpret_cast(detail::curve::detail::d_variables[cv]), + reinterpret_cast(detail::curve::detail::d_variables[cv] + (I * J * K * W * sizeof(T)))); // Read-write flag resides in 8 bits at the end of the buffer + } + return ReadOnlyDeviceMacroProperty(nullptr, nullptr); +#else + return ReadOnlyDeviceMacroProperty(reinterpret_cast(detail::curve::detail::d_variables[cv])); +#endif +} +template +__device__ __forceinline__ DeviceMacroProperty DeviceEnvironment::getMacroProperty(const char(&name)[N]) const { + detail::curve::Curve::VariableHash cvh = MACRO_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); + const auto cv = detail::curve::Curve::getVariable(cvh); +#if !defined(SEATBELTS) || SEATBELTS + if (cv == detail::curve::Curve::UNKNOWN_VARIABLE) { + DTHROW("Environment macro property name: %s was not found.\n", name); + } else if (detail::curve::detail::d_sizes[cv] != sizeof(T)) { + DTHROW("Environment macro property with name: %s type size mismatch %llu != %llu.\n", name, detail::curve::detail::d_sizes[cv], sizeof(T)); + } else if (detail::curve::detail::d_lengths[cv] != I * J * K * W) { + DTHROW("Environment macro property with name: %s total length mismatch (%u != %u).\n", name, detail::curve::detail::d_lengths[cv], I * J * K * W); + } else { + return DeviceMacroProperty(reinterpret_cast(detail::curve::detail::d_variables[cv]), + reinterpret_cast(detail::curve::detail::d_variables[cv] + (I * J * K * W * sizeof(T)))); // Read-write flag resides in 8 bits at the end of the buffer + } + return DeviceMacroProperty(nullptr, nullptr); +#else + return DeviceMacroProperty(reinterpret_cast(detail::curve::detail::d_variables[cv])); +#endif +} #endif // __CUDACC_RTC__ } // namespace flamegpu diff --git a/include/flamegpu/runtime/utility/DeviceMacroProperty.cuh b/include/flamegpu/runtime/utility/DeviceMacroProperty.cuh new file mode 100644 index 000000000..efb5a16ee --- /dev/null +++ b/include/flamegpu/runtime/utility/DeviceMacroProperty.cuh @@ -0,0 +1,452 @@ +#ifndef INCLUDE_FLAMEGPU_RUNTIME_UTILITY_DEVICEMACROPROPERTY_CUH_ +#define INCLUDE_FLAMEGPU_RUNTIME_UTILITY_DEVICEMACROPROPERTY_CUH_ + +#include +#include +#include + +namespace flamegpu { + +/** + * This template class is used for conveniently converting a multi-dimensional pointer to an array + * Theoretically the compiler should be able to optimise away most of it at compile time + */ +template +class ReadOnlyDeviceMacroProperty { + protected: + T* ptr; +#if !defined(SEATBELTS) || SEATBELTS + /** + * This flag is used by seatbelts to check for a read/atomic-write conflict + * Reading sets 1<<0 + * Writing sets 1<<1 + */ + unsigned int* read_write_flag; + /** + * Utility function for setting/checking read flag + */ + __device__ void setCheckReadFlag() const; + /** + * Utility function for setting/checking write flag + */ + __device__ void setCheckWriteFlag() const; +#endif + + public: +#if !defined(SEATBELTS) || SEATBELTS + /** + * Constructor + * @param _ptr Pointer to buffer + * @param _rwf Pointer to read_write_flag + */ + __device__ explicit ReadOnlyDeviceMacroProperty(T* _ptr, unsigned int* _rwf); +#else + /** + * Constructor + * @param _ptr Pointer to buffer + */ + __device__ explicit ReadOnlyDeviceMacroProperty(T* _ptr); +#endif + /** + * Access the next dimension of the array + * @throws exception::DeviceError If i >= I. + * @throws exception::DeviceError If template arguments I, J, K , W are all 1. Which denotes the macro property has no dimensions remaining to be indexed. + */ + __device__ __forceinline__ ReadOnlyDeviceMacroProperty operator[](unsigned int i) const; + /** + * Read-only access to the current element + * @throws exception::DeviceError If template arguments I, J, K , W are not all 1. Which denotes the macro property has dimensions remaining to be indexed. + */ + __device__ __forceinline__ operator T() const; +}; +template +class DeviceMacroProperty : public ReadOnlyDeviceMacroProperty { + public: +#if !defined(SEATBELTS) || SEATBELTS + /** + * Constructor + * @param _ptr Pointer to buffer + * @param _rwf Pointer to read_write_flag + */ + __device__ explicit DeviceMacroProperty(T* _ptr, unsigned int *_rwf); +#else + /** + * Constructor + * @param _ptr Pointer to buffer + */ + __device__ explicit DeviceMacroProperty(T* _ptr); +#endif + /** + * Access the next dimension of the array + * @throws exception::DeviceError If i >= I. + * @throws exception::DeviceError If template arguments I, J, K , W are all 1. Which denotes the macro property has no dimensions remaining to be indexed. + */ + __device__ __forceinline__ DeviceMacroProperty operator[](unsigned int i) const; + /** + * atomic add + * @param val The 2nd operand + * @return a reference to this + * Note, taking value of the returned object will fail, due to the risk of atomic conflicts + * @note Only suitable where T is type int32_t, uint32_t, uint64_t, float, double + */ + __device__ __forceinline__ DeviceMacroProperty& operator +=(const T& val); + /** + * atomic subtraction + * @param val The 2nd operand + * @return a reference to this + * Note, taking value of the returned object will fail, due to the risk of atomic conflicts + * @note Only suitable where T is type int32_t or uint32_t + */ + __device__ __forceinline__ DeviceMacroProperty& operator -=(const T& val); + /** + * atomic add + * @param val The 2nd operand + * @return (this + val) + * @note Only suitable where T is type int32_t, uint32_t, uint64_t, float, double + */ + __device__ __forceinline__ T operator+(const T& val) const; + /** + * atomic subtraction + * @param val The 2nd operand + * @return (this - val) + * @note Only suitable where T is type int32_t or uint32_t + */ + __device__ __forceinline__ T operator-(const T& val) const; + /** + * atomic pre-increment + * @return the value after the increment operation is performed + * @note A reference is not returned, as this makes it easy to fall into the trap (race condition) of reading it's value. + * @note Only suitable where T is type uint32_t + */ + __device__ __forceinline__ T operator++(); + /** + * atomic pre-decrement + * @return the value after the decrement operation is performed + * @note A reference is not returned, as this makes it easy to fall into the trap (race condition) of reading it's value. + * @note Only suitable where T is type uint32_t + */ + __device__ __forceinline__ T operator--(); + /** + * atomic post-increment + * @return the value before increment operation is performed + * @note Only suitable where T is type uint32_t + */ + __device__ __forceinline__ T operator++(int); + /** + * atomic post-decrement + * @return the value before increment operation is performed + * @note Only suitable where T is type uint32_t + */ + __device__ __forceinline__ T operator--(int); + /** + * atomic min + * @return min(this, val) + * @note Only suitable where T is type int32_t, uint32_t, uint64_t + */ + __device__ __forceinline__ T min(T val); + /** + * atomic max + * @return max(this, val) + * @note Only suitable where T is type int32_t, uint32_t, uint64_t + */ + __device__ __forceinline__ T max(T val); + /** + * atomic compare and swap + * Computes (old == compare ? val : old) + * @return old + * @note Only suitable where T is type int32_t, uint32_t, uint64_t, uint16_t + */ + __device__ __forceinline__ T CAS(T compare, T val); + /** + * atomic exchange + * Returns the current value stored in the element, and replaces it with val + * @return the value before the exchange + * @note Only suitable for 32/64 bit types + */ + __device__ __forceinline__ T exchange(T val); +}; + +#if !defined(SEATBELTS) || SEATBELTS +template +__device__ __forceinline__ ReadOnlyDeviceMacroProperty::ReadOnlyDeviceMacroProperty(T* _ptr, unsigned int* _rwf) + : ptr(_ptr) + , read_write_flag(_rwf) +{ } +template +__device__ __forceinline__ DeviceMacroProperty::DeviceMacroProperty(T* _ptr, unsigned int* _rwf) + : ReadOnlyDeviceMacroProperty(_ptr, _rwf) +{ } +#ifdef __CUDACC__ +template +__device__ void ReadOnlyDeviceMacroProperty::setCheckReadFlag() const { + const unsigned int old = atomicOr(read_write_flag, 1u << 0); + if (old & 1u << 1) { + DTHROW("DeviceMacroProperty read and atomic write operations cannot be mixed in the same layer, as this may cause race conditions.\n"); + } +} +template +__device__ void ReadOnlyDeviceMacroProperty::setCheckWriteFlag() const { + const unsigned int old = atomicOr(read_write_flag, 1u << 1); + if (old & 1u << 0) { + DTHROW("DeviceMacroProperty read and atomic write operations cannot be mixed in the same layer as this may cause race conditions.\n"); + } +} +#endif +#else +template +__device__ __forceinline__ ReadOnlyDeviceMacroProperty::ReadOnlyDeviceMacroProperty(T* _ptr) + :ptr(_ptr) +{ } +template +__device__ __forceinline__ DeviceMacroProperty::DeviceMacroProperty(T* _ptr) + : ReadOnlyDeviceMacroProperty(_ptr) +{ } +#endif +template +__device__ __forceinline__ ReadOnlyDeviceMacroProperty ReadOnlyDeviceMacroProperty::operator[](unsigned int i) const { +#if !defined(SEATBELTS) || SEATBELTS + if (I == 1 && J == 1 && K == 1 && W == 1) { + DTHROW("Indexing error, property has less dimensions.\n"); + } else if (i >= I) { + DTHROW("Indexing error, out of bounds %u >= %u.\n", i, I); + } else if (this->ptr == nullptr) { + return ReadOnlyDeviceMacroProperty(nullptr, nullptr); + } +#endif + // (i * J * K * W) + (j * K * W) + (k * W) + w +#if !defined(SEATBELTS) || SEATBELTS + return ReadOnlyDeviceMacroProperty(this->ptr + (i * J * K * W), this->read_write_flag); +#else + return DeviceMacroProperty(this->ptr + (i * J * K * W)); +#endif +} +template +__device__ __forceinline__ DeviceMacroProperty DeviceMacroProperty::operator[](unsigned int i) const { +#if !defined(SEATBELTS) || SEATBELTS + if (I == 1 && J == 1 && K == 1 && W == 1) { + DTHROW("Indexing error, property has less dimensions.\n"); + } else if (i >= I) { + DTHROW("Indexing error, out of bounds %u >= %u.\n", i, I); + } else if (this->ptr == nullptr) { + return DeviceMacroProperty(nullptr, nullptr); + } +#endif + // (i * J * K * W) + (j * K * W) + (k * W) + w +#if !defined(SEATBELTS) || SEATBELTS + return DeviceMacroProperty(this->ptr + (i * J * K * W), this->read_write_flag); +#else + return DeviceMacroProperty(this->ptr + (i * J * K * W)); +#endif +} +template +__device__ __forceinline__ ReadOnlyDeviceMacroProperty::operator T() const { +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckReadFlag(); +#endif + return *this->ptr; +} +template +__device__ __forceinline__ DeviceMacroProperty& DeviceMacroProperty::operator+=(const T& val) { + static_assert(std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value, "atomic add only supports the types int32_t/uint32_t/uint64_t/float/double."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return *this; + } + this->setCheckWriteFlag(); +#endif + atomicAdd(this->ptr, val); + return *this; +} +template +__device__ __forceinline__ DeviceMacroProperty& DeviceMacroProperty::operator-=(const T& val) { + static_assert(std::is_same::value || std::is_same::value, "atomic subtract only supports the types int32_t/uint32_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return *this; + } + this->setCheckWriteFlag(); +#endif + atomicSub(this->ptr, val); + return *this; +} +template +__device__ __forceinline__ T DeviceMacroProperty::operator+(const T& val) const { + static_assert(std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value, "atomic add only supports the types int32_t/uint32_t/uint64_t/float/double."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return atomicAdd(this->ptr, val) + val; +} +template +__device__ __forceinline__ T DeviceMacroProperty::operator-(const T& val) const { + static_assert(std::is_same::value || std::is_same::value, "atomic subtract only supports the types int32_t/uint32_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return atomicSub(this->ptr, val) - val; +} +template +__device__ __forceinline__ T DeviceMacroProperty::operator++() { + static_assert(std::is_same::value, "atomic increment only supports the type uint32_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return *this; + } + this->setCheckWriteFlag(); +#endif + const T old = atomicInc(this->ptr, std::numeric_limits::max()); + return ((old >= std::numeric_limits::max()) ? 0 : (old + 1)); +} + +template +__device__ __forceinline__ T DeviceMacroProperty::operator--() { + static_assert(std::is_same::value, "atomic decrement only supports the type uint32_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return *this; + } + this->setCheckWriteFlag(); +#endif + const T old = atomicDec(this->ptr, std::numeric_limits::max()); + return (((old == 0) || (old > std::numeric_limits::max())) ? std::numeric_limits::max() : (old - 1)); +} +template +__device__ __forceinline__ T DeviceMacroProperty::operator++(int) { + static_assert(std::is_same::value, "atomic increment only supports the type uint32_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return atomicInc(this->ptr, std::numeric_limits::max()); +} + +template +__device__ __forceinline__ T DeviceMacroProperty::operator--(int) { + static_assert(std::is_same::value, "atomic decrement only supports the type uint32_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return atomicDec(this->ptr, std::numeric_limits::max()); +} +template +__device__ __forceinline__ T DeviceMacroProperty::min(T val) { + static_assert(std::is_same::value || + std::is_same::value || + std::is_same::value, "atomic min only supports the types int32_t/uint32_t/uint64_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return std::min(atomicMin(this->ptr, val), val); +} +template +__device__ __forceinline__ T DeviceMacroProperty::max(T val) { + static_assert(std::is_same::value || + std::is_same::value || + std::is_same::value, "atomic max only supports the types int32_t/uint32_t/uint64_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return std::max(atomicMax(this->ptr, val), val); +} +template +__device__ __forceinline__ T DeviceMacroProperty::CAS(T compare, T val) { + static_assert(std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value, "atomic compare and swap only supports the types int32_t/uint32_t/uint64_t/uint16_t."); +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + return atomicCAS(this->ptr, compare, val); +} + +// GCC doesn't like seeing atomicExch with host compiler +#ifdef __CUDACC__ +#pragma diag_suppress = initialization_not_reachable +template +__device__ __forceinline__ T DeviceMacroProperty::exchange(T val) { + static_assert(std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value || + std::is_same::value, "atomic exchange only supports the types int32_t/int64_t/uint32_t/uint64_t/float/double."); + static_assert(sizeof(uint64_t) == sizeof(unsigned long long int), "uint64_t != unsigned long long int."); // NOLINT(runtime/int) +#if !defined(SEATBELTS) || SEATBELTS + if (I != 1 || J != 1 || K != 1 || W != 1) { + DTHROW("Indexing error, property has more dimensions.\n"); + } else if (this->ptr == nullptr) { + return { }; + } + this->setCheckWriteFlag(); +#endif + if (sizeof(T) == sizeof(uint64_t)) { // Convert all 64 bit types to unsigned long long int (can't build as uint64_t on gcc) + const unsigned long long int rval = atomicExch(reinterpret_cast(this->ptr), *reinterpret_cast(&val)); // NOLINT(runtime/int) + return *reinterpret_cast(&rval); + } + // else 32-bit + const uint32_t rval = atomicExch(reinterpret_cast(this->ptr), *reinterpret_cast(&val)); + return *reinterpret_cast(&rval); + // return atomicExch(this->ptr, val); +} +#pragma diag_default = initialization_not_reachable +#endif + +} // namespace flamegpu + +#endif // INCLUDE_FLAMEGPU_RUNTIME_UTILITY_DEVICEMACROPROPERTY_CUH_ diff --git a/include/flamegpu/runtime/utility/HostEnvironment.cuh b/include/flamegpu/runtime/utility/HostEnvironment.cuh index c2856edb1..ae8375cdd 100644 --- a/include/flamegpu/runtime/utility/HostEnvironment.cuh +++ b/include/flamegpu/runtime/utility/HostEnvironment.cuh @@ -11,8 +11,10 @@ #include #include +#include "flamegpu/gpu/CUDAMacroEnvironment.h" #include "flamegpu/gpu/detail/CUDAErrorChecking.cuh" #include "flamegpu/runtime/utility/EnvironmentManager.cuh" +#include "flamegpu/runtime/utility/HostMacroProperty.cuh" namespace flamegpu { @@ -33,11 +35,15 @@ class HostEnvironment { /** * Constructor, to be called by HostAPI */ - explicit HostEnvironment(const unsigned int &instance_id); + explicit HostEnvironment(const unsigned int &instance_id, CUDAMacroEnvironment &_macro_env); /** * Provides access to EnvironmentManager singleton */ EnvironmentManager &env_mgr; + /** + * Provides access to macro properties for the instance + */ + CUDAMacroEnvironment& macro_env; /** * Access to instance id of the CUDASimulation * This is used to augment all variable names @@ -131,6 +137,20 @@ class HostEnvironment { */ template std::vector setPropertyArray(const std::string &name, const std::vector &value) const; +#endif + /** + * Returns an interface for accessing the named host macro property + * @param name The name of the environment macro property to return + */ + template + HostMacroProperty getMacroProperty(const std::string& name) const; +#ifdef SWIG + /** + * None-templated dimensions version of getMacroProperty() for SWIG interface + * @param name The name of the environment macro property to return + */ + template + HostMacroProperty_swig getMacroProperty_swig(const std::string& name) const; #endif }; @@ -194,6 +214,17 @@ std::vector HostEnvironment::getPropertyArray(const std::string& name) const } #endif // SWIG +template +HostMacroProperty HostEnvironment::getMacroProperty(const std::string& name) const { + return macro_env.getProperty(name); +} + +#ifdef SWIG +template +HostMacroProperty_swig HostEnvironment::getMacroProperty_swig(const std::string& name) const { + return macro_env.getProperty_swig(name); +} +#endif } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_RUNTIME_UTILITY_HOSTENVIRONMENT_CUH_ diff --git a/include/flamegpu/runtime/utility/HostMacroProperty.cuh b/include/flamegpu/runtime/utility/HostMacroProperty.cuh new file mode 100644 index 000000000..c1a271d68 --- /dev/null +++ b/include/flamegpu/runtime/utility/HostMacroProperty.cuh @@ -0,0 +1,572 @@ +#ifndef INCLUDE_FLAMEGPU_RUNTIME_UTILITY_HOSTMACROPROPERTY_CUH_ +#define INCLUDE_FLAMEGPU_RUNTIME_UTILITY_HOSTMACROPROPERTY_CUH_ + +#include +#include +#include +#include +#include +#include + +namespace flamegpu { + +struct HostMacroProperty_MetaData { + /** + * Constructor + */ + HostMacroProperty_MetaData(void* _d_base_ptr, const std::array& _dims, size_t _type_size, + bool _device_read_flag, const std::string &name) + : h_base_ptr(nullptr) + , d_base_ptr(static_cast(_d_base_ptr)) + , dims(_dims) + , elements(dims[0] * dims[1] * dims[2] * dims[3]) + , type_size(_type_size) + , has_changed(false) + , device_read_flag(_device_read_flag) + , property_name(name) + { } + ~HostMacroProperty_MetaData() { + upload(); + if (h_base_ptr) + std::free(h_base_ptr); + } + /** + * Download data + */ + void download() { + if (!h_base_ptr) { + h_base_ptr = static_cast(malloc(elements * type_size)); + gpuErrchk(cudaMemcpy(h_base_ptr, d_base_ptr, elements * type_size, cudaMemcpyDeviceToHost)); + has_changed = false; + } + } + /** + * Upload data + */ + void upload() { + if (h_base_ptr && has_changed) { +#if !defined(SEATBELTS) || SEATBELTS + if (device_read_flag) { + THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, " + "in HostMacroProperty_MetaData::upload()\n", + property_name.c_str()); + } +#endif + gpuErrchk(cudaMemcpy(d_base_ptr, h_base_ptr, elements * type_size, cudaMemcpyHostToDevice)); + has_changed = false; + } + } + char* h_base_ptr; + char* d_base_ptr; + std::array dims; + unsigned int elements; + size_t type_size; + bool has_changed; + bool device_read_flag; + std::string property_name; +}; + +/** + * This template class is used for conveniently converting a multi-dimensional pointer to an array + * Theoretically the compiler should be able to optimise away most of it at compile time + */ +template +class HostMacroProperty { + /** + * Data ptr to current section + */ + ptrdiff_t offset; + std::shared_ptr metadata; + + public: + /** + * Root Constructor + * @param _metadata Metadata struct + */ + explicit HostMacroProperty(const std::shared_ptr& _metadata); + /** + * Sub Constructor + * @param _metadata Preallocated metadata struct, shared with other current instances of same host macro property + * @param _offset Pointer to buffer in device memory + */ + HostMacroProperty(const std::shared_ptr &_metadata, const ptrdiff_t& _offset); + + /** + * Access the next dimension of the array + * @throws exception::OutOfBoundsException If i >= I. + * @throws exception::InvalidOperation If template arguments I, J, K , W are all 1. Which denotes the macro property has no dimensions remaining to be indexed. + */ + HostMacroProperty operator[](unsigned int i) const; + /** + * Read/Write access to the current element + * @throws exception::InvalidOperation If template arguments I, J, K , W are not all 1. Which denotes the macro property has more dimensions remaining to be indexed. + */ + operator T(); + /** + * Read-only access to the current element + * @throws exception::InvalidOperation If template arguments I, J, K , W are not all 1. Which denotes the macro property has more dimensions remaining to be indexed. + */ + operator T() const; + /** + * Assign value + * @param val New value of the element + * @throws exception::InvalidOperation If template arguments I, J, K , W are not all 1. Which denotes the macro property has more dimensions remaining to be indexed. + */ + HostMacroProperty& operator=(const T &val); + /** + * Zero's the selected area of the array + */ + void zero(); + + /** + * Increment operators + */ + HostMacroProperty& operator++(); + T operator++(int); + HostMacroProperty& operator--(); + T operator--(int); + +// Swig breaks if it see's auto return type +#ifndef SWIG + /** + * Arithmetic operators + */ + template + auto operator+(const T2& b) const; + template + auto operator-(const T2& b) const; + template + auto operator*(const T2& b) const; + template + auto operator/(const T2& b) const; + template + auto operator%(const T2& b) const; +#endif + /** + * Assignment operators + */ + // HostMacroProperty& operator=(const T& b); // Defined above + template + HostMacroProperty& operator+=(const T2& b); + template + HostMacroProperty& operator-=(const T2& b); + template + HostMacroProperty& operator*=(const T2& b); + template + HostMacroProperty& operator/=(const T2& b); + template + HostMacroProperty& operator%=(const T2& b); + + private: + /** + * Validate and download data, set changed flag + */ + T& _get(); + /** + * Validate and download data + */ + T& _get() const; +}; + +#ifdef SWIG +/** + * This template class is used for conveniently converting a multi-dimensional pointer to an array + * Theoretically the compiler should be able to optimise away most of it at compile time + */ +template +class HostMacroProperty_swig { + /** + * This exists in place of the template args found in HostMacroProperty + */ + const std::array dimensions; + /** + * Data ptr to current section + */ + ptrdiff_t offset; + std::shared_ptr metadata; + + public: + /** + * Root Constructor + * @param _metadata Metadata struct + */ + explicit HostMacroProperty_swig(const std::shared_ptr &_metadata); + /** + * Sub Constructor + * @param _metadata Preallocated metadata struct, shared with other current instances of same host macro property + * @param _offset Pointer to buffer in device memory + * @param _dimensions In place of HostMacroProperty's template args + */ + HostMacroProperty_swig(const std::shared_ptr& _metadata, const ptrdiff_t& _offset, const std::array& _dimensions); + /** + * Access the next dimension of the array + * @throws exception::OutOfBoundsException If i >= I. + * @throws exception::InvalidOperation If the macro property has no dimensions remaining to be indexed. + */ + HostMacroProperty_swig __getitem__(unsigned int i) const; + /** + * Set the item at position in the array + * @note This is due to the fact Python doesn't less us override operator= directly. + * @throws exception::OutOfBoundsException If i >= I. + * @throws exception::InvalidOperation If the macro property has no dimensions remaining to be indexed. + */ + void __setitem__(unsigned int i, const T &val); + /** + * Explicit set method, as we lack operator= in python + */ + void set(T val); + + int __int__(); + int64_t __long__(); + double __float__(); + bool __bool__(); + + bool __eq__(const T& other) const; + bool __ne__(const T &other) const; + bool __lt__(const T& other) const; + bool __le__(const T& other) const; + bool __gt__(const T& other) const; + bool __ge__(const T& other) const; + + /** + * Required for iterable + */ + unsigned int __len__() const { return dimensions[0]; } + + /** + * Required for python extension + */ + T get() const; + /** + * Zero's the selected area of the array + */ + void zero(); +}; +#endif + +template +HostMacroProperty::HostMacroProperty(const std::shared_ptr &_metadata) +: offset(0u) +, metadata(_metadata) +{ } +template +HostMacroProperty::HostMacroProperty(const std::shared_ptr &_metadata, const ptrdiff_t & _offset) + : offset(_offset) + , metadata(_metadata) +{ } +template +HostMacroProperty HostMacroProperty::operator[](unsigned int i) const { + if (I == 1 && J == 1 && K == 1 && W == 1) { + THROW exception::InvalidOperation("Indexing error, property has less dimensions.\n"); + } else if (i >= I) { + THROW exception::OutOfBoundsException("Indexing error, out of bounds %u >= %u.\n", i, I); + } + // (i * J * K * W) + (j * K * W) + (k * W) + w + return HostMacroProperty(metadata, offset + (i * J * K * W)); +} +template +HostMacroProperty::operator T() { + if (I != 1 || J != 1 || K != 1 || W != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + // Must assume changed + metadata->has_changed = true; + return *(reinterpret_cast(metadata->h_base_ptr) + offset); +} +template +HostMacroProperty::operator T() const { + metadata->download(); + return *(reinterpret_cast(metadata->h_base_ptr) + offset); +} + +template +void HostMacroProperty::zero() { + if (metadata->h_base_ptr) { + // Memset on host + memset(reinterpret_cast(metadata->h_base_ptr) + offset, 0, I * J * K * W * metadata->type_size); + metadata->has_changed = true; + } else { +#if !defined(SEATBELTS) || SEATBELTS + if (metadata->device_read_flag) { + THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, " + "in HostMacroProperty::zero()\n", + metadata->property_name.c_str()); + } +#endif + // Memset on device + gpuErrchk(cudaMemset(reinterpret_cast(metadata->d_base_ptr) + offset, 0, I * J * K * W * metadata->type_size)); + } +} +template +T& HostMacroProperty::_get() const { + if (I != 1 || J != 1 || K != 1 || W != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset]; +} +template +T& HostMacroProperty::_get() { + if (I != 1 || J != 1 || K != 1 || W != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + metadata->has_changed = true; + return reinterpret_cast(metadata->h_base_ptr)[offset]; +} + + +template +HostMacroProperty& HostMacroProperty::operator++() { + ++_get(); + return *this; +} +template +T HostMacroProperty::operator++(int) { + T &t = _get(); + T ret = t; + ++t; + return ret; +} +template +HostMacroProperty& HostMacroProperty::operator--() { + --_get(); + return *this; +} +template +T HostMacroProperty::operator--(int) { + T& t = _get(); + T ret = t; + --t; + return ret; +} + +#ifndef SWIG +template +template +auto HostMacroProperty::operator+(const T2& b) const { + return _get() + b; +} +template +template +auto HostMacroProperty::operator-(const T2& b) const { + return _get() - b; +} +template +template +auto HostMacroProperty::operator*(const T2& b) const { + return _get() * b; +} +template +template +auto HostMacroProperty::operator/(const T2& b) const { + return _get() / b; +} +template +template +auto HostMacroProperty::operator%(const T2& b) const { + return _get() % b; +} +#endif + +template +template +HostMacroProperty& HostMacroProperty::operator+=(const T2& b) { + _get() += b; + return *this; +} +template +template +HostMacroProperty& HostMacroProperty::operator-=(const T2& b) { + _get() -= b; + return *this; +} +template +template +HostMacroProperty& HostMacroProperty::operator*=(const T2& b) { + _get() *= b; + return *this; +} +template +template +HostMacroProperty& HostMacroProperty::operator/=(const T2& b) { + _get() /= b; + return *this; +} +template +template +HostMacroProperty& HostMacroProperty::operator%=(const T2& b) { + _get() %= b; + return *this; +} + +template +HostMacroProperty& HostMacroProperty::operator=(const T& val) { + if (I != 1 || J != 1 || K != 1 || W != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + reinterpret_cast(metadata->h_base_ptr)[offset] = val; + metadata->has_changed = true; + return *this; +} + +// SWIG versions +#ifdef SWIG +template +HostMacroProperty_swig::HostMacroProperty_swig(const std::shared_ptr& _metadata) + : dimensions(_metadata->dims) + , offset(0u) + , metadata(_metadata) +{ } +template +HostMacroProperty_swig::HostMacroProperty_swig(const std::shared_ptr& _metadata, const ptrdiff_t& _offset, const std::array &_dimensions) + : dimensions(_dimensions) + , offset(_offset) + , metadata(_metadata) +{ } +template +HostMacroProperty_swig HostMacroProperty_swig::__getitem__(unsigned int i) const { + if (dimensions[0] == 1 && dimensions[1] == 1 && dimensions[2] == 1 && dimensions[3] == 1) { + THROW exception::InvalidOperation("Indexing error, property has less dimensions.\n"); + } else if (i >= dimensions[0]) { + THROW exception::OutOfBoundsException("Indexing error, out of bounds %u >= %u.\n", i, dimensions[0]); + } + // (i * J * K * W) + (j * K * W) + (k * W) + w + return HostMacroProperty_swig(metadata, offset + (i * dimensions[1] * dimensions[2] * dimensions[3]), { dimensions[1], dimensions[2], dimensions[3], 1 }); +} + +template +void HostMacroProperty_swig::zero() { + if (metadata->h_base_ptr) { + // Memset on host + memset(reinterpret_cast(metadata->h_base_ptr) + offset, 0, dimensions[0] * dimensions[1] * dimensions[2] * dimensions[3] * metadata->type_size); + metadata->has_changed = true; + } else { + // Memset on device + gpuErrchk(cudaMemset(reinterpret_cast(metadata->d_base_ptr) + offset, 0, dimensions[0] * dimensions[1] * dimensions[2] * dimensions[3] * metadata->type_size)); + } +} +template +void HostMacroProperty_swig::set(T val) { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + reinterpret_cast(metadata->h_base_ptr)[offset] = val; + metadata->has_changed = true; +} +template +void HostMacroProperty_swig::__setitem__(unsigned int i, const T& val) { + if (dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } else if (i >= dimensions[0]) { + THROW exception::InvalidOperation("Indexing out of bounds %u >= %u.\n", i, dimensions[0]); + } + metadata->download(); + unsigned int t_offset = offset + (i * dimensions[1] * dimensions[2] * dimensions[3]); + reinterpret_cast(metadata->h_base_ptr)[t_offset] = val; + metadata->has_changed = true; +} +template +int HostMacroProperty_swig::__int__() { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return static_cast(reinterpret_cast(metadata->h_base_ptr)[offset]); +} +template +int64_t HostMacroProperty_swig::__long__() { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return static_cast(reinterpret_cast(metadata->h_base_ptr)[offset]); +} +template +double HostMacroProperty_swig::__float__() { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return static_cast(reinterpret_cast(metadata->h_base_ptr)[offset]); +} +template +bool HostMacroProperty_swig::__bool__() { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return static_cast(reinterpret_cast(metadata->h_base_ptr)[offset]); +} +template +bool HostMacroProperty_swig::__eq__(const T& other) const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset] == other; +} +template +bool HostMacroProperty_swig::__ne__(const T& other) const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset] != other; +} +template +bool HostMacroProperty_swig::__lt__(const T& other) const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset] < other; +} +template +bool HostMacroProperty_swig::__le__(const T& other) const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset] <= other; +} +template +bool HostMacroProperty_swig::__gt__(const T& other) const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset] > other; +} +template +bool HostMacroProperty_swig::__ge__(const T& other) const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset] >= other; +} +// template +// T HostMacroProperty_swig::__mod__(const T& other) { +// if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { +// THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); +// } +// metadata->download(); +// return reinterpret_cast(metadata->h_base_ptr)[offset] % other; +// } +template +T HostMacroProperty_swig::get() const { + if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) { + THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n"); + } + metadata->download(); + return reinterpret_cast(metadata->h_base_ptr)[offset]; +} +#endif + +} // namespace flamegpu + +#endif // INCLUDE_FLAMEGPU_RUNTIME_UTILITY_HOSTMACROPROPERTY_CUH_ diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index fd761c49a..85a7a67f4 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -132,6 +132,7 @@ SET(SRC_INCLUDE ${FLAMEGPU_ROOT}/include/flamegpu/gpu/CUDAFatAgent.h ${FLAMEGPU_ROOT}/include/flamegpu/gpu/CUDAFatAgentStateList.h ${FLAMEGPU_ROOT}/include/flamegpu/gpu/CUDAScatter.cuh + ${FLAMEGPU_ROOT}/include/flamegpu/gpu/CUDAMacroEnvironment.h ${FLAMEGPU_ROOT}/include/flamegpu/sim/AgentInterface.h ${FLAMEGPU_ROOT}/include/flamegpu/sim/AgentLoggingConfig.h ${FLAMEGPU_ROOT}/include/flamegpu/sim/AgentLoggingConfig_Reductions.cuh @@ -183,8 +184,10 @@ SET(SRC_INCLUDE ${FLAMEGPU_ROOT}/include/flamegpu/runtime/messaging/MessageBucket/MessageBucketDevice.cuh ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/AgentRandom.cuh ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/DeviceEnvironment.cuh + ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/DeviceMacroProperty.cuh ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/EnvironmentManager.cuh ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/HostEnvironment.cuh + ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/HostMacroProperty.cuh ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/HostRandom.cuh ${FLAMEGPU_ROOT}/include/flamegpu/runtime/utility/RandomManager.cuh ${FLAMEGPU_ROOT}/include/flamegpu/util/Any.h @@ -237,6 +240,7 @@ SET(SRC_FLAMEGPU ${FLAMEGPU_ROOT}/src/flamegpu/gpu/CUDAScatter.cu ${FLAMEGPU_ROOT}/src/flamegpu/gpu/CUDASimulation.cu ${FLAMEGPU_ROOT}/src/flamegpu/gpu/CUDAEnsemble.cu + ${FLAMEGPU_ROOT}/src/flamegpu/gpu/CUDAMacroEnvironment.cu ${FLAMEGPU_ROOT}/src/flamegpu/sim/AgentLoggingConfig.cu ${FLAMEGPU_ROOT}/src/flamegpu/sim/LoggingConfig.cu ${FLAMEGPU_ROOT}/src/flamegpu/sim/LogFrame.cu diff --git a/src/flamegpu/gpu/CUDAAgent.cu b/src/flamegpu/gpu/CUDAAgent.cu index 658dd99e0..083dcd966 100644 --- a/src/flamegpu/gpu/CUDAAgent.cu +++ b/src/flamegpu/gpu/CUDAAgent.cu @@ -517,7 +517,7 @@ void CUDAAgent::clearFunctionCondition(const std::string &state) { fat_agent->setConditionState(fat_index, state, 0); } -void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool function_condition) { +void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, const CUDAMacroEnvironment ¯o_env, bool function_condition) { // Generate the dynamic curve header detail::curve::CurveRTCHost &curve_header = *rtc_header_map.emplace(function_condition ? func.name + "_condition" : func.name, std::make_unique()).first->second; @@ -581,6 +581,9 @@ void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool f } } + // Set Environment macro properties in curve + macro_env.mapRTCVariables(curve_header); + std::string header_filename = std::string(func.rtc_func_name).append("_impl"); if (function_condition) header_filename.append("_condition"); diff --git a/src/flamegpu/gpu/CUDAMacroEnvironment.cu b/src/flamegpu/gpu/CUDAMacroEnvironment.cu new file mode 100644 index 000000000..39e37da68 --- /dev/null +++ b/src/flamegpu/gpu/CUDAMacroEnvironment.cu @@ -0,0 +1,198 @@ +#include "flamegpu/gpu/CUDAMacroEnvironment.h" + +#include "flamegpu/model/EnvironmentDescription.h" +#include "flamegpu/gpu/CUDASimulation.h" +#include "flamegpu/model/AgentFunctionData.cuh" +#include "flamegpu/model/SubEnvironmentData.h" +#include "flamegpu/runtime/detail/curve/curve_rtc.cuh" + +namespace flamegpu { + +const char CUDAMacroEnvironment::MACRO_NAMESPACE_STRING[18] = "MACRO_ENVIRONMENT"; + +CUDAMacroEnvironment::CUDAMacroEnvironment(const EnvironmentDescription& description, const CUDASimulation& _cudaSimulation) + : MACRO_NAMESPACE_HASH(detail::curve::Curve::variableRuntimeHash(MACRO_NAMESPACE_STRING)) + , cudaSimulation(_cudaSimulation) { + assert(MACRO_NAMESPACE_HASH == DeviceEnvironment::MACRO_NAMESPACE_HASH()); // Host and Device namespace const's do not match + for (const auto &p : description.getMacroPropertiesMap()) { + properties.emplace(p.first, MacroEnvProp(p.second.type, p.second.type_size, p.second.elements)); + } +} + +void CUDAMacroEnvironment::init() { + for (auto &prop : properties) { + if (!prop.second.d_ptr) { + size_t buffer_size = prop.second.type_size + * prop.second.elements[0] + * prop.second.elements[1] + * prop.second.elements[2] + * prop.second.elements[3]; +#if !defined(SEATBELTS) || SEATBELTS + buffer_size += sizeof(unsigned int); // Extra uint is used as read-write flag by seatbelts +#endif + gpuErrchk(cudaMalloc(&prop.second.d_ptr, buffer_size)); + gpuErrchk(cudaMemset(prop.second.d_ptr, 0, buffer_size)); + } + } + mapRuntimeVariables(); +} + +void CUDAMacroEnvironment::init(const SubEnvironmentData& mapping, const CUDAMacroEnvironment &master_macro_env) { + // Map local properties + for (auto& prop : properties) { + if (!prop.second.d_ptr) { + auto sub = mapping.macro_properties.find(prop.first); + if (sub == mapping.macro_properties.end()) { + // If it's a local macro property + size_t buffer_size = prop.second.type_size + * prop.second.elements[0] + * prop.second.elements[1] + * prop.second.elements[2] + * prop.second.elements[3]; +#if !defined(SEATBELTS) || SEATBELTS + buffer_size += sizeof(unsigned int); // Extra uint is used as read-write flag by seatbelts +#endif + gpuErrchk(cudaMalloc(&prop.second.d_ptr, buffer_size)); + gpuErrchk(cudaMemset(prop.second.d_ptr, 0, buffer_size)); + } else { + // If it's a mapped sub macro property + auto mmp = master_macro_env.properties.find(sub->second); + if (mmp != master_macro_env.properties.end() + && mmp->second.d_ptr + && mmp->second.elements == prop.second.elements + && mmp->second.type == prop.second.type) { + prop.second.d_ptr = mmp->second.d_ptr; + prop.second.is_sub = true; + } else { + THROW exception::UnknownInternalError("Unable to initialise mapped macro property '%s' to '%s', this should not have failed, " + "in CUDAMacroEnvironment::init()\n", + prop.first.c_str(), sub->second.c_str()); + } + } + } + } + // Pass them all to CURVE + mapRuntimeVariables(); +} +void CUDAMacroEnvironment::free() { + unmapRuntimeVariables(); + for (auto& prop : properties) { + if (prop.second.d_ptr) { + if (!prop.second.is_sub) { + gpuErrchk(cudaFree(prop.second.d_ptr)); + } + prop.second.d_ptr = nullptr; + } + } +} +void CUDAMacroEnvironment::purge() { + for (auto& prop : properties) + prop.second.d_ptr = nullptr; +} + +void CUDAMacroEnvironment::mapRuntimeVariables() const { + auto& curve = detail::curve::Curve::getInstance(); + // loop through the agents variables to map each variable name using cuRVE + for (const auto& mmp : properties) { + // map using curve + const detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); + + // get the agent variable size + const unsigned int length = mmp.second.elements[0] * mmp.second.elements[1] * mmp.second.elements[2] * mmp.second.elements[3]; + +#ifdef _DEBUG + const detail::curve::Curve::Variable cv = curve.registerVariableByHash(var_hash + MACRO_NAMESPACE_HASH + cudaSimulation.getInstanceID(), mmp.second.d_ptr, mmp.second.type_size, length); + if (cv != static_cast((var_hash + MACRO_NAMESPACE_HASH + cudaSimulation.getInstanceID()) % detail::curve::Curve::MAX_VARIABLES)) { + fprintf(stderr, "detail::curve::Curve Warning: Environment macro property '%s' has a collision and may work improperly.\n", mmp.first.c_str()); + } +#else + curve.registerVariableByHash(var_hash + MACRO_NAMESPACE_HASH + cudaSimulation.getInstanceID(), mmp.second.d_ptr, mmp.second.type_size, length); +#endif + } +} + +void CUDAMacroEnvironment::unmapRuntimeVariables() const { + // loop through the agents variables to unmap each property using cuRVE + for (const auto& mmp : properties) { + const detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); + detail::curve::Curve::getInstance().unregisterVariableByHash(var_hash + MACRO_NAMESPACE_HASH + cudaSimulation.getInstanceID()); + } +} +void CUDAMacroEnvironment::mapRTCVariables(detail::curve::CurveRTCHost& curve_header) const { + for (const auto &p : properties) { + curve_header.registerEnvMacroProperty(p.first.c_str(), p.second.d_ptr, p.second.type.name(), p.second.type_size, p.second.elements); + } +} +void CUDAMacroEnvironment::unmapRTCVariables(detail::curve::CurveRTCHost& curve_header) const { + for (const auto &p : properties) { + curve_header.unregisterEnvMacroProperty(p.first.c_str()); + } +} +#if !defined(SEATBELTS) || SEATBELTS +void CUDAMacroEnvironment::resetFlagsAsync(const std::vector &streams) { + unsigned int i = 0; + for (const auto& prop : properties) { + if (prop.second.d_ptr) { + const size_t buffer_size = prop.second.type_size + * prop.second.elements[0] + * prop.second.elements[1] + * prop.second.elements[2] + * prop.second.elements[3]; + gpuErrchk(cudaMemsetAsync(static_cast(prop.second.d_ptr) + buffer_size, 0 , sizeof(unsigned int), streams[i++%streams.size()])); + } + } + // Disable the sync here, users must sync themselves + // if (properties.size()) { + // gpuErrchk(cudaDeviceSynchronize()); + // } +} +bool CUDAMacroEnvironment::getDeviceReadFlag(const std::string& property_name) { + const auto prop = properties.find(property_name); + if (prop == properties.end()) { + THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, " + "in CUDAMacroEnvironment::getDeviceReadFlag()\n", + property_name.c_str()); + } + const size_t buffer_size = prop->second.type_size + * prop->second.elements[0] + * prop->second.elements[1] + * prop->second.elements[2] + * prop->second.elements[3]; + unsigned int ret = 0; + gpuErrchk(cudaMemcpy(&ret, static_cast(prop->second.d_ptr) + buffer_size, sizeof(unsigned int), cudaMemcpyDeviceToHost)); + return (ret & 1u << 0); +} +bool CUDAMacroEnvironment::getDeviceWriteFlag(const std::string& property_name) { + const auto prop = properties.find(property_name); + if (prop == properties.end()) { + THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, " + "in CUDAMacroEnvironment::getDeviceReadFlag()\n", + property_name.c_str()); + } + const size_t buffer_size = prop->second.type_size + * prop->second.elements[0] + * prop->second.elements[1] + * prop->second.elements[2] + * prop->second.elements[3]; + unsigned int ret = 0; + gpuErrchk(cudaMemcpy(&ret, static_cast(prop->second.d_ptr) + buffer_size, sizeof(unsigned int), cudaMemcpyDeviceToHost)); + return (ret & 1u << 1); +} +unsigned int CUDAMacroEnvironment::getDeviceRWFlags(const std::string& property_name) { + const auto prop = properties.find(property_name); + if (prop == properties.end()) { + THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, " + "in CUDAMacroEnvironment::getDeviceReadFlag()\n", + property_name.c_str()); + } + const size_t buffer_size = prop->second.type_size + * prop->second.elements[0] + * prop->second.elements[1] + * prop->second.elements[2] + * prop->second.elements[3]; + unsigned int ret = 0; + gpuErrchk(cudaMemcpy(&ret, static_cast(prop->second.d_ptr) + buffer_size, sizeof(unsigned int), cudaMemcpyDeviceToHost)); + return ret; +} +#endif +} // namespace flamegpu diff --git a/src/flamegpu/gpu/CUDASimulation.cu b/src/flamegpu/gpu/CUDASimulation.cu index 959244e51..be54102a3 100644 --- a/src/flamegpu/gpu/CUDASimulation.cu +++ b/src/flamegpu/gpu/CUDASimulation.cu @@ -62,6 +62,7 @@ CUDASimulation::CUDASimulation(const std::shared_ptr &_model) , elapsedMillisecondsInitFunctions(0.f) , elapsedMillisecondsExitFunctions(0.f) , elapsedMillisecondsRTCInitialisation(0.f) + , macro_env(*_model->environment, *this) , run_log(std::make_unique()) , streams(std::vector()) , singletons(nullptr) @@ -114,6 +115,7 @@ bool CUDASimulation::detectPureRTC(const std::shared_ptr& _mode CUDASimulation::CUDASimulation(const std::shared_ptr &submodel_desc, CUDASimulation *master_model) : Simulation(submodel_desc, master_model) , step_count(0) + , macro_env(*submodel_desc->submodel->environment, *this) , run_log(std::make_unique()) , streams(std::vector()) , singletons(nullptr) @@ -201,6 +203,7 @@ CUDASimulation::~CUDASimulation() { message_map.clear(); submodel_map.clear(); host_api.reset(); + macro_env.free(); #ifdef VISUALISATION visualisation.reset(); #endif @@ -834,6 +837,12 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un // Execute the host functions. layerHostFunctions(layer, layerIndex); +#if !defined(SEATBELTS) || SEATBELTS + // Reset macro-environment read-write flags + // Note this does not synchronise threads, it relies on synchronizeAllStreams() post host fns + macro_env.resetFlagsAsync(streams); +#endif + // Synchronise after the host layer functions to ensure that the device is up to date? This can potentially be removed. this->synchronizeAllStreams(); } @@ -1292,6 +1301,7 @@ void CUDASimulation::initialiseSingletons() { singletons->scatter.purge(); } EnvironmentManager::getInstance().purge(); + macro_env.purge(); // Reset flag DEVICE_HAS_RESET_CHECK = 0; // Any value that doesnt match DEVICE_HAS_RESET_FLAG gpuErrchk(cudaMemcpyToSymbol(DEVICE_HAS_RESET, &DEVICE_HAS_RESET_CHECK, sizeof(unsigned int))); @@ -1307,7 +1317,7 @@ void CUDASimulation::initialiseSingletons() { singletons->rng.reseed(getSimulationConfig().random_seed); // Pass created RandomManager to host api - host_api = std::make_unique(*this, singletons->rng, singletons->scatter, agentOffsets, agentData, 0, getStream(0)); // Host fns are currently all serial + host_api = std::make_unique(*this, singletons->rng, singletons->scatter, agentOffsets, agentData, macro_env, 0, getStream(0)); // Host fns are currently all serial for (auto &cm : message_map) { cm.second->init(singletons->scatter, 0); @@ -1316,8 +1326,10 @@ void CUDASimulation::initialiseSingletons() { // Populate the environment properties if (!submodel) { singletons->environment.init(instance_id, *model->environment, isPureRTC); + macro_env.init(); } else { singletons->environment.init(instance_id, *model->environment, isPureRTC, mastermodel->getInstanceID(), *submodel->subenvironment); + macro_env.init(*submodel->subenvironment, mastermodel->macro_env); } // Propagate singleton init to submodels @@ -1364,12 +1376,12 @@ void CUDASimulation::initialiseRTC() { // check rtc source to see if this is a RTC function if (!it_f->second->rtc_source.empty()) { // create CUDA agent RTC function by calling addInstantitateRTCFunction on CUDAAgent with AgentFunctionData - a_it->second->addInstantitateRTCFunction(*it_f->second); + a_it->second->addInstantitateRTCFunction(*it_f->second, macro_env); } // check rtc source to see if the function condition is an rtc condition if (!it_f->second->rtc_condition_source.empty()) { // create CUDA agent RTC function condition by calling addInstantitateRTCFunction on CUDAAgent with AgentFunctionData - a_it->second->addInstantitateRTCFunction(*it_f->second, true); + a_it->second->addInstantitateRTCFunction(*it_f->second, macro_env, true); } } } diff --git a/src/flamegpu/model/EnvironmentDescription.cu b/src/flamegpu/model/EnvironmentDescription.cu index 9568735b9..35d5d9e4e 100644 --- a/src/flamegpu/model/EnvironmentDescription.cu +++ b/src/flamegpu/model/EnvironmentDescription.cu @@ -17,20 +17,18 @@ bool EnvironmentDescription::operator==(const EnvironmentDescription& rhs) const auto _v = rhs.properties.find(v.first); if (_v == rhs.properties.end()) return false; - if (v.second.isConst != _v->second.isConst - || v.second.data.elements != _v->second.data.elements - || v.second.data.type != _v->second.data.type) + if (v.second != _v->second) + return false; + } + return true; + } + if (macro_properties.size() == rhs.macro_properties.size()) { + for (auto& v : macro_properties) { + auto _v = rhs.macro_properties.find(v.first); + if (_v == rhs.macro_properties.end()) + return false; + if (v.second != _v->second) return false; - if (v.second.data.ptr == _v->second.data.ptr && - v.second.data.length == _v->second.data.length) - continue; - if (v.second.data.length == _v->second.data.length) { - for (size_t i = 0; i < v.second.data.length; ++i) - if (reinterpret_cast(v.second.data.ptr)[i] != reinterpret_cast(_v->second.data.ptr)[i]) - return false; - continue; - } - return false; } return true; } @@ -58,5 +56,8 @@ bool EnvironmentDescription::getConst(const std::string &name) { const std::unordered_map EnvironmentDescription::getPropertiesMap() const { return properties; } +const std::unordered_map EnvironmentDescription::getMacroPropertiesMap() const { + return macro_properties; +} } // namespace flamegpu diff --git a/src/flamegpu/model/LayerDescription.cpp b/src/flamegpu/model/LayerDescription.cpp index 4fead2896..8cdea7ddb 100644 --- a/src/flamegpu/model/LayerDescription.cpp +++ b/src/flamegpu/model/LayerDescription.cpp @@ -137,7 +137,7 @@ void LayerDescription::addHostFunction(FLAMEGPU_HOST_FUNCTION_POINTER func_p) { "in LayerDescription::addHostFunction()\n"); } if (!layer->host_functions.empty() || !layer->agent_functions.empty() || !layer->host_functions_callbacks.empty()) { - THROW exception::InvalidLayerMember("A layer containing agent functions or a host function may not also contain a host function, " + THROW exception::InvalidLayerMember("A layer containing agent functions or a host function may not have another host function added, " "in LayerDescription::addHostFunction()\n"); } if (!layer->host_functions.insert(func_p).second) { diff --git a/src/flamegpu/model/SubEnvironmentData.cpp b/src/flamegpu/model/SubEnvironmentData.cpp index 9e96bc32c..e7f0dcf8e 100644 --- a/src/flamegpu/model/SubEnvironmentData.cpp +++ b/src/flamegpu/model/SubEnvironmentData.cpp @@ -10,7 +10,7 @@ bool SubEnvironmentData::operator==(const SubEnvironmentData& rhs) const { if (this == &rhs) // They point to same object return true; // Compare members - if (properties == rhs.properties) { + if (properties == rhs.properties && macro_properties == rhs.macro_properties) { return true; } return false; @@ -25,6 +25,7 @@ SubEnvironmentData::SubEnvironmentData(const std::shared_ptr &m , parent(_parent) , description(model ? new SubEnvironmentDescription(model, this) : nullptr) { properties.insert(other.properties.begin(), other.properties.end()); + macro_properties.insert(other.macro_properties.begin(), other.macro_properties.end()); } SubEnvironmentData::SubEnvironmentData( const std::shared_ptr &model, diff --git a/src/flamegpu/model/SubEnvironmentDescription.cpp b/src/flamegpu/model/SubEnvironmentDescription.cpp index 1e8fbb688..ea4154168 100644 --- a/src/flamegpu/model/SubEnvironmentDescription.cpp +++ b/src/flamegpu/model/SubEnvironmentDescription.cpp @@ -63,28 +63,90 @@ void SubEnvironmentDescription::mapProperty(const std::string &sub_property_name // Properties match, create mapping data->properties.emplace(sub_property_name, master_property_name); } +void SubEnvironmentDescription::mapMacroProperty(const std::string& sub_property_name, const std::string& master_property_name) { + // Sub macro property exists + auto subEnv = data->subEnvironment.lock(); + if (!subEnv) { + THROW exception::InvalidParent("SubEnvironment pointer has expired, " + "in SubEnvironmentDescription::mapMacroProperty()\n"); + } + const auto subProp = subEnv->macro_properties.find(sub_property_name); + if (subProp == subEnv->macro_properties.end()) { + const auto parent = data->parent.lock(); + THROW exception::InvalidEnvProperty("SubModel '%s's Environment does not contain macro property '%s', " + "in SubEnvironmentDescription::mapMacroProperty()\n", parent ? parent->submodel->name.c_str() : "?", sub_property_name.c_str()); + } + // Master macro property exists + auto masterEnv = data->masterEnvironment.lock(); + if (!masterEnv) { + THROW exception::InvalidParent("MasterEnvironment pointer has expired, " + "in SubEnvironmentDescription::mapMacroProperty()\n"); + } + const auto masterProp = masterEnv->macro_properties.find(master_property_name); + if (masterProp == masterEnv->macro_properties.end()) { + THROW exception::InvalidEnvProperty("MasterEnvironment does not contain macro property '%s', " + "in SubEnvironmentDescription::mapMacroProperty()\n", master_property_name.c_str()); + } + // Sub macro property has not been bound yet + if (data->macro_properties.find(sub_property_name) != data->macro_properties.end()) { + const auto parent = data->parent.lock(); + THROW exception::InvalidEnvProperty("SubModel '%s's Environment macro property '%s' has already been mapped, " + "in SubEnvironmentDescription::mapMacroProperty()\n", parent ? parent->submodel->name.c_str() : "?", sub_property_name.c_str()); + } + // Master macro property has already been bound + for (auto& v : data->macro_properties) { + if (v.second == master_property_name) { + THROW exception::InvalidEnvProperty("MasterEnvironment macro property '%s' has already been mapped, " + "in SubEnvironmentDescription::mapMacroProperty()\n", master_property_name.c_str()); + } + } + // Check macro properties are the same + if (subProp->second.type != masterProp->second.type) { + THROW exception::InvalidEnvProperty("Macro property types do not match, '%s' != '%s', " + "in SubEnvironmentDescription::mapMacroProperty()\n", subProp->second.type.name(), masterProp->second.type.name()); + } + if (subProp->second.elements != masterProp->second.elements) { + THROW exception::InvalidEnvProperty("Macro property dimensions do not match, (%u, %u, %u, %u) != (%u, %u, %u, %u)", + "in SubEnvironmentDescription::mapMacroProperty()\n", + subProp->second.elements[0], subProp->second.elements[1], subProp->second.elements[2], subProp->second.elements[3], + masterProp->second.elements[0], masterProp->second.elements[1], masterProp->second.elements[2], masterProp->second.elements[3]); + } + // Macro properties match, create mapping + data->macro_properties.emplace(sub_property_name, master_property_name); +} std::string SubEnvironmentDescription::getPropertyMapping(const std::string &sub_property_name) { const auto v = data->properties.find(sub_property_name); if (v != data->properties.end()) return v->second; THROW exception::InvalidAgentState("SubEnvironment property '%s', either does not exist or has not been mapped yet, " - "in SubAgentDSubEnvironmentDescriptionescription::getPropertyMapping()\n", sub_property_name.c_str()); + "in SubEnvironmentDescription::getPropertyMapping()\n", sub_property_name.c_str()); +} +std::string SubEnvironmentDescription::getMacroPropertyMapping(const std::string& sub_property_name) { + const auto v = data->macro_properties.find(sub_property_name); + if (v != data->properties.end()) + return v->second; + THROW exception::InvalidAgentState("SubEnvironment macro property '%s', either does not exist or has not been mapped yet, " + "in SubEnvironmentDescription::getMacroPropertyMapping()\n", sub_property_name.c_str()); } +void SubEnvironmentDescription::autoMap() { + autoMapProperties(); + autoMapMacroProperties(); +} void SubEnvironmentDescription::autoMapProperties() { // Sub property exists auto subEnv = data->subEnvironment.lock(); if (!subEnv) { THROW exception::InvalidParent("SubEnvironment pointer has expired, " - "in SubEnvironmentDescription::mapProperty()\n"); + "in SubEnvironmentDescription::autoMapProperties()\n"); } // Master property exists auto masterEnv = data->masterEnvironment.lock(); if (!masterEnv) { THROW exception::InvalidParent("MasterEnvironment pointer has expired, " - "in SubEnvironmentDescription::mapProperty()\n"); + "in SubEnvironmentDescription::autoMapProperties()\n"); } for (auto &subProp : subEnv->properties) { auto masterProp = masterEnv->properties.find(subProp.first); @@ -99,5 +161,30 @@ void SubEnvironmentDescription::autoMapProperties() { } } } +void SubEnvironmentDescription::autoMapMacroProperties() { + // Sub property exists + auto subEnv = data->subEnvironment.lock(); + if (!subEnv) { + THROW exception::InvalidParent("SubEnvironment pointer has expired, " + "in SubEnvironmentDescription::autoMapMacroProperties()\n"); + } + // Master property exists + auto masterEnv = data->masterEnvironment.lock(); + if (!masterEnv) { + THROW exception::InvalidParent("MasterEnvironment pointer has expired, " + "in SubEnvironmentDescription::autoMapMacroProperties()\n"); + } + for (auto& subProp : subEnv->macro_properties) { + auto masterProp = masterEnv->macro_properties.find(subProp.first); + // If there exists variable with same name in both environments + if (masterProp != masterEnv->macro_properties.end()) { + // Check properties are the same + if ((subProp.second.type == masterProp->second.type) && + (subProp.second.elements == masterProp->second.elements)) { + data->macro_properties.emplace(subProp.first, masterProp->first); // Doesn't actually matter, both strings are equal + } + } + } +} } // namespace flamegpu diff --git a/src/flamegpu/runtime/HostAPI.cu b/src/flamegpu/runtime/HostAPI.cu index 98c071665..21387707c 100644 --- a/src/flamegpu/runtime/HostAPI.cu +++ b/src/flamegpu/runtime/HostAPI.cu @@ -12,10 +12,11 @@ HostAPI::HostAPI(CUDASimulation &_agentModel, CUDAScatter &_scatter, const AgentOffsetMap &_agentOffsets, AgentDataMap &_agentData, + CUDAMacroEnvironment ¯o_env, const unsigned int& _streamId, cudaStream_t _stream) : random(rng) - , environment(_agentModel.getInstanceID()) + , environment(_agentModel.getInstanceID(), macro_env) , agentModel(_agentModel) , d_cub_temp(nullptr) , d_cub_temp_size(0) diff --git a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp index 91351b7b3..9b1df2a0c 100644 --- a/src/flamegpu/runtime/detail/curve/curve_rtc.cpp +++ b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp @@ -172,24 +172,35 @@ __device__ __forceinline__ void Curve::setNewAgentArrayVariable(const char(&name // has to be included after definition of curve namespace #include "flamegpu/runtime/utility/DeviceEnvironment.cuh" +//#include "flamegpu/runtime/utility/DeviceMacroProperty.cuh" namespace flamegpu { template -__device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N]) const { +__device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[N]) const { $DYNAMIC_ENV_GETVARIABLE_IMPL } template -__device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N], const unsigned int &index) const { +__device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[N], const unsigned int &index) const { $DYNAMIC_ENV_GETARRAYVARIABLE_IMPL } template -__device__ __forceinline__ bool DeviceEnvironment::containsProperty(const char(&name)[N]) const { +__device__ __forceinline__ bool ReadOnlyDeviceEnvironment::containsProperty(const char(&name)[N]) const { $DYNAMIC_ENV_CONTAINTS_IMPL } + +template +__device__ __forceinline__ ReadOnlyDeviceMacroProperty ReadOnlyDeviceEnvironment::getMacroProperty(const char(&name)[N]) const { +$DYNAMIC_ENV_GETREADONLYMACROPROPERTY_IMPL +} +template +__device__ __forceinline__ DeviceMacroProperty DeviceEnvironment::getMacroProperty(const char(&name)[N]) const { +$DYNAMIC_ENV_GETMACROPROPERTY_IMPL +} + } // namespace flamegpu #endif // CURVE_RTC_DYNAMIC_H_ @@ -330,6 +341,26 @@ void CurveRTCHost::unregisterEnvVariable(const char* propertyName) { THROW exception::UnknownInternalError("Environment property '%s' not found when removing environment property, in CurveRTCHost::unregisterEnvVariable()", propertyName); } } +void CurveRTCHost::registerEnvMacroProperty(const char* propertyName, void *d_ptr, const char* type, size_t type_size, const std::array &dimensions) { + RTCEnvMacroPropertyProperties props; + props.type = CurveRTCHost::demangle(type); + props.dimensions = dimensions; + props.d_ptr = d_ptr; + props.h_data_ptr = nullptr; + props.type_size = type_size; + if (!RTCEnvMacroProperties.emplace(propertyName, props).second) { + THROW exception::UnknownInternalError("Environment property with name '%s' is already registered, in CurveRTCHost::registerEnvMacroProperty()", propertyName); + } +} + +void CurveRTCHost::unregisterEnvMacroProperty(const char* propertyName) { + auto i = RTCEnvMacroProperties.find(propertyName); + if (i != RTCEnvMacroProperties.end()) { + RTCEnvMacroProperties.erase(propertyName); + } else { + THROW exception::UnknownInternalError("Environment macro property '%s' not found when removing environment property, in CurveRTCHost::unregisterEnvMacroProperty()", propertyName); + } +} void CurveRTCHost::initHeaderEnvironment() { @@ -343,6 +374,7 @@ void CurveRTCHost::initHeaderEnvironment() { messageOut_data_offset = data_buffer_size; data_buffer_size += messageOut_variables.size() * sizeof(void*); messageIn_data_offset = data_buffer_size; data_buffer_size += messageIn_variables.size() * sizeof(void*); newAgent_data_offset = data_buffer_size; data_buffer_size += newAgent_variables.size() * sizeof(void*); + envMacro_data_offset = data_buffer_size; data_buffer_size += RTCEnvMacroProperties.size() * sizeof(void*); variables << "__constant__ char " << getVariableSymbolName() << "[" << data_buffer_size << "];\n"; setHeaderPlaceholder("$DYNAMIC_VARIABLES", variables.str()); // generate Environment::get func implementation ($DYNAMIC_ENV_GETVARIABLE_IMPL) @@ -419,6 +451,76 @@ void CurveRTCHost::initHeaderEnvironment() { containsEnvVariableImpl << " return false;\n"; setHeaderPlaceholder("$DYNAMIC_ENV_CONTAINTS_IMPL", containsEnvVariableImpl.str()); } + // generate Environment::getMacroProperty func implementation ($DYNAMIC_ENV_GETREADONLYMACROPROPERTY_IMPL) + { + size_t ct = 0; + std::stringstream getMacroPropertyImpl; + for (std::pair element : RTCEnvMacroProperties) { + RTCEnvMacroPropertyProperties props = element.second; + getMacroPropertyImpl << " if (strings_equal(name, \"" << element.first << "\")) {\n"; + getMacroPropertyImpl << "#if !defined(SEATBELTS) || SEATBELTS\n"; + getMacroPropertyImpl << " if(sizeof(T) != " << element.second.type_size << ") {\n"; + getMacroPropertyImpl << " DTHROW(\"Environment macro property '%s' type mismatch.\\n\", name);\n"; + getMacroPropertyImpl << " } else if (I != " << element.second.dimensions[0] << " ||\n"; + getMacroPropertyImpl << " J != " << element.second.dimensions[1] << " |\n"; + getMacroPropertyImpl << " K != " << element.second.dimensions[2] << " |\n"; + getMacroPropertyImpl << " W != " << element.second.dimensions[3] << ") {\n"; + getMacroPropertyImpl << " DTHROW(\"Environment macro property '%s' dimensions do not match (%u, %u, %u, %u) != (%u, %u, %u, %u).\\n\", name,\n"; + getMacroPropertyImpl << " I, J, K, W, " << element.second.dimensions[0] << ", " << element.second.dimensions[1] << ", " << element.second.dimensions[2] << ", " << element.second.dimensions[3] << ");\n"; + getMacroPropertyImpl << " } else {\n"; + getMacroPropertyImpl << " return ReadOnlyDeviceMacroProperty(*reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << envMacro_data_offset + (ct * sizeof(void*)) << "),\n"; + // Read-write flag resides in 8 bits at the end of the buffer + getMacroPropertyImpl << " reinterpret_cast(*reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << envMacro_data_offset + (ct * sizeof(void*)) << ") + (I * J * K * W * sizeof(T))));\n"; + getMacroPropertyImpl << " }\n"; + getMacroPropertyImpl << "#else\n"; + getMacroPropertyImpl << " return ReadOnlyDeviceMacroProperty(*reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << envMacro_data_offset + (ct * sizeof(void*)) << "));\n"; + getMacroPropertyImpl << "#endif\n"; + getMacroPropertyImpl << " };\n"; + ++ct; + } + getMacroPropertyImpl << "#if !defined(SEATBELTS) || SEATBELTS\n"; + getMacroPropertyImpl << " DTHROW(\"Environment macro property '%s' was not found.\\n\", name);\n"; + getMacroPropertyImpl << " return ReadOnlyDeviceMacroProperty(nullptr, nullptr);\n"; + getMacroPropertyImpl << "#else\n"; + getMacroPropertyImpl << " return ReadOnlyDeviceMacroProperty(nullptr);\n"; + getMacroPropertyImpl << "#endif\n"; + setHeaderPlaceholder("$DYNAMIC_ENV_GETREADONLYMACROPROPERTY_IMPL", getMacroPropertyImpl.str()); + } + // generate Environment::getMacroProperty func implementation ($DYNAMIC_ENV_GETMACROPROPERTY_IMPL) + { + size_t ct = 0; + std::stringstream getMacroPropertyImpl; + for (std::pair element : RTCEnvMacroProperties) { + RTCEnvMacroPropertyProperties props = element.second; + getMacroPropertyImpl << " if (strings_equal(name, \"" << element.first << "\")) {\n"; + getMacroPropertyImpl << "#if !defined(SEATBELTS) || SEATBELTS\n"; + getMacroPropertyImpl << " if(sizeof(T) != " << element.second.type_size << ") {\n"; + getMacroPropertyImpl << " DTHROW(\"Environment macro property '%s' type mismatch.\\n\", name);\n"; + getMacroPropertyImpl << " } else if (I != " << element.second.dimensions[0] << " ||\n"; + getMacroPropertyImpl << " J != " << element.second.dimensions[1] << " |\n"; + getMacroPropertyImpl << " K != " << element.second.dimensions[2] << " |\n"; + getMacroPropertyImpl << " W != " << element.second.dimensions[3] << ") {\n"; + getMacroPropertyImpl << " DTHROW(\"Environment macro property '%s' dimensions do not match (%u, %u, %u, %u) != (%u, %u, %u, %u).\\n\", name,\n"; + getMacroPropertyImpl << " I, J, K, W, " << element.second.dimensions[0] << ", " << element.second.dimensions[1] << ", " << element.second.dimensions[2] << ", " << element.second.dimensions[3] << ");\n"; + getMacroPropertyImpl << " } else {\n"; + getMacroPropertyImpl << " return DeviceMacroProperty(*reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << envMacro_data_offset + (ct * sizeof(void*)) << "),\n"; + // Read-write flag resides in 8 bits at the end of the buffer + getMacroPropertyImpl << " reinterpret_cast(*reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << envMacro_data_offset + (ct * sizeof(void*)) << ") + (I * J * K * W * sizeof(T))));\n"; + getMacroPropertyImpl << " }\n"; + getMacroPropertyImpl << "#else\n"; + getMacroPropertyImpl << " return DeviceMacroProperty(*reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << envMacro_data_offset + (ct * sizeof(void*)) << "));\n"; + getMacroPropertyImpl << "#endif\n"; + getMacroPropertyImpl << " };\n"; + ++ct; + } + getMacroPropertyImpl << "#if !defined(SEATBELTS) || SEATBELTS\n"; + getMacroPropertyImpl << " DTHROW(\"Environment macro property '%s' was not found.\\n\", name);\n"; + getMacroPropertyImpl << " return DeviceMacroProperty(nullptr, nullptr);\n"; + getMacroPropertyImpl << "#else\n"; + getMacroPropertyImpl << " return DeviceMacroProperty(nullptr);\n"; + getMacroPropertyImpl << "#endif\n"; + setHeaderPlaceholder("$DYNAMIC_ENV_GETMACROPROPERTY_IMPL", getMacroPropertyImpl.str()); + } } void CurveRTCHost::initHeaderSetters() { // generate setAgentVariable func implementation ($DYNAMIC_SETAGENTVARIABLE_IMPL) @@ -919,6 +1021,12 @@ void CurveRTCHost::initDataBuffer() { for (auto &element : newAgent_variables) { element.second.h_data_ptr = h_data_buffer + newAgent_data_offset + (ct++ * sizeof(void*)); } + ct = 0; + for (auto& element : RTCEnvMacroProperties) { + element.second.h_data_ptr = h_data_buffer + envMacro_data_offset + (ct++ * sizeof(void*)); + // Env macro properties don't update, so fill them as we go + memcpy(element.second.h_data_ptr, &element.second.d_ptr, sizeof(void*)); + } } void CurveRTCHost::setFileName(const std::string &filename) { diff --git a/src/flamegpu/runtime/utility/HostEnvironment.cu b/src/flamegpu/runtime/utility/HostEnvironment.cu index f63f3a7a3..f5d0ded29 100644 --- a/src/flamegpu/runtime/utility/HostEnvironment.cu +++ b/src/flamegpu/runtime/utility/HostEnvironment.cu @@ -2,8 +2,9 @@ namespace flamegpu { -HostEnvironment::HostEnvironment(const unsigned int &_instance_id) +HostEnvironment::HostEnvironment(const unsigned int &_instance_id, CUDAMacroEnvironment& _macro_env) : env_mgr(EnvironmentManager::getInstance()) + , macro_env(_macro_env) , instance_id(_instance_id) { } } // namespace flamegpu diff --git a/swig/python/flamegpu.i b/swig/python/flamegpu.i index 5a0f54c7b..b8ae3bb2e 100644 --- a/swig/python/flamegpu.i +++ b/swig/python/flamegpu.i @@ -12,6 +12,8 @@ #pragma SWIG nowarn=362 // operator++ ignored (not supported by python) #pragma SWIG nowarn=383 +// operator-- ignored (not supported by python) +#pragma SWIG nowarn=384 // Warning 451 Setting a const char * variable may leak memory. Fix is to use a std::string instead? #pragma SWIG nowarn=451 // LoggingConfig.h:65: Function flamegpu::LoggingConfig::Any(ReductionFn) must have a return type. Ignored. This actually a typedef function pointer? @@ -226,6 +228,7 @@ class FLAMEGPUIterator(object): raise StopIteration %} + // Exception handling. /** Exception handling * FLAMEGPURuntimeException class is a wrapper class to replace specific instances of FLAMEGPUException. It is constructed with a error mesage and type which can be queried to give the original error and the original exception class. @@ -524,6 +527,9 @@ class ModelVis; // Must wrap these prior to HostAPI where they are used to avoid issues with no default constructors etc. %include "flamegpu/runtime/utility/HostRandom.cuh" + +%nodefaultctor flamegpu::HostMacroProperty_swig; +%include "flamegpu/runtime/utility/HostMacroProperty.cuh" %include "flamegpu/runtime/utility/HostEnvironment.cuh" %include "flamegpu/runtime/HostNewAgentAPI.h" @@ -657,6 +663,72 @@ TEMPLATE_VARIABLE_INSTANTIATE_ID(getProperty, flamegpu::HostEnvironment::getProp TEMPLATE_VARIABLE_INSTANTIATE_ID(getPropertyArray, flamegpu::HostEnvironment::getPropertyArray) TEMPLATE_VARIABLE_INSTANTIATE_ID(setProperty, flamegpu::HostEnvironment::setProperty) TEMPLATE_VARIABLE_INSTANTIATE_ID(setPropertyArray, flamegpu::HostEnvironment::setPropertyArray) +TEMPLATE_VARIABLE_INSTANTIATE_ID(getMacroProperty, flamegpu::HostEnvironment::getMacroProperty_swig) + +// Instance template versions of the HostMacroProperty class +// Extend HostMacroProperty so that it is python iterable +%extend flamegpu::HostMacroProperty_swig { + %pythoncode { + def __iter__(self): + return FLAMEGPUIterator(self) + def __add__(self, other): + return self.get() + other; + def __radd__(self, other): + return other + self.get(); + def __iadd__(self, other): + self.set(self.get() + other); + return self; + def __sub__(self, other): + return self.get() - other; + def __rsub__(self, other): + return other - self.get(); + def __isub__(self, other): + self.set(self.get() - other); + return self; + def __mul__(self, other): + return self.get() * other; + def __rmul__(self, other): + return other * self.get(); + def __imul__(self, other): + self.set(self.get() * other); + return self; + def __pow__(self, other): + return self.get() ** other; + def __rpow__(self, other): + return other ** self.get(); + def __ipow__(self, other): + self.set(self.get() ** other); + return self; + def __truediv__(self, other): + return self.get() / other; + def __rtruediv__(self, other): + return other / self.get(); + def __itruediv__(self, other): + try: + self.set(self.get() / other); + except: + raise FLAMEGPURuntimeException("__itruediv__ does not support the used type combination as it would lead to type conversion of the host object.", "unsupported type") + return self; + def __floordiv__(self, other): + return self.get() // other; + def __rfloordiv__(self, other): + return other // self.get(); + def __ifloordiv__(self, other): + self.set(self.get() // other); + return self; + def __mod__(self, other): + return self.get() % other; + def __rmod__(self, other): + return other % self.get(); + def __imod__(self, other): + try: + self.set(self.get() % other); + except: + raise FLAMEGPURuntimeException("__imod__ does not support the used type combination as it would lead to type conversion of the host object.", "unsupported type") + return self; + } +} +TEMPLATE_VARIABLE_INSTANTIATE_ID(HostMacroProperty, flamegpu::HostMacroProperty_swig) // Instantiate template versions of host agent functions from the API TEMPLATE_VARIABLE_INSTANTIATE_ID(getVariable, flamegpu::HostNewAgentAPI::getVariable) @@ -668,6 +740,7 @@ TEMPLATE_VARIABLE_INSTANTIATE_ID(setVariableArray, flamegpu::HostNewAgentAPI::se // Instantiate template versions of environment description functions from the API TEMPLATE_VARIABLE_INSTANTIATE_ID(newProperty, flamegpu::EnvironmentDescription::newProperty) TEMPLATE_VARIABLE_INSTANTIATE_ID(newPropertyArray, flamegpu::EnvironmentDescription::newPropertyArray) +TEMPLATE_VARIABLE_INSTANTIATE_ID(newMacroProperty, flamegpu::EnvironmentDescription::newMacroProperty_swig) TEMPLATE_VARIABLE_INSTANTIATE_ID(getProperty, flamegpu::EnvironmentDescription::getProperty) TEMPLATE_VARIABLE_INSTANTIATE_ID(getPropertyArray, flamegpu::EnvironmentDescription::getPropertyArray) //TEMPLATE_VARIABLE_INSTANTIATE_ID(getPropertyAt, flamegpu::EnvironmentDescription::getPropertyArrayAtIndex) @@ -738,8 +811,6 @@ TEMPLATE_VARIABLE_INSTANTIATE_ID(newVariableArray, flamegpu::MessageArray2D::Des TEMPLATE_VARIABLE_INSTANTIATE_ID(newVariableArray, flamegpu::MessageArray3D::Description::newVariableArray) TEMPLATE_VARIABLE_INSTANTIATE_ID(newVariableArray, flamegpu::MessageBucket::Description::newVariableArray) -// Instantiate template versions of host random functions from the API - // Instantiate template versions of host random functions from the API TEMPLATE_VARIABLE_INSTANTIATE_FLOATS(uniform, flamegpu::HostRandom::uniformNoRange) TEMPLATE_VARIABLE_INSTANTIATE_INTS(uniform, flamegpu::HostRandom::uniformRange) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index afae36d0d..d7d87729d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -16,6 +16,7 @@ SET(TESTS_SRC ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/gpu/test_cuda_ensemble.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/gpu/test_gpu_validation.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/gpu/test_cuda_subagent.cu + ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/gpu/test_cuda_submacroenvironment.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/io/test_io.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/io/test_logging.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/io/test_logging_exceptions.cu @@ -34,17 +35,20 @@ SET(TESTS_SRC ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/sim/test_host_functions.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/sim/test_RunPlan.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/sim/test_RunPlanVector.cu - ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_agent_environment.cu + ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_environment.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_agent_function_conditions.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_agent_random.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_agent_state_transition.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_agent_creation.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_api.cu + ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_environment.cu + ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_device_macro_property.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_environment_manager.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_api.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_agent_sort.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_agent_creation.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_environment.cu + ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_macro_property.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/test_host_random.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/host_reduction/test_min.cu ${CMAKE_CURRENT_SOURCE_DIR}/test_cases/runtime/host_reduction/test_max.cu diff --git a/tests/swig/python/gpu/test_cuda_submacroenvironment.py b/tests/swig/python/gpu/test_cuda_submacroenvironment.py new file mode 100644 index 000000000..4051b67bb --- /dev/null +++ b/tests/swig/python/gpu/test_cuda_submacroenvironment.py @@ -0,0 +1,169 @@ +""" + Tests of submodel functionality of class: CUDAMacroEnvironment + ReadTest: Test that HostMacroProperty can read data (written via agent fn's) + WriteTest: Test that HostMacroProperty can write data (read via agent fn's) + ZeroTest: Test that HostMacroProperty can zero data (read via agent fn's) +""" +import pytest +from unittest import TestCase +from pyflamegpu import * + +class ExitAlways(pyflamegpu.HostFunctionConditionCallback): + def run(self, FLAMEGPU): + return pyflamegpu.EXIT; + +class Host_Write_5(pyflamegpu.HostFunctionCallback): + def run(self, FLAMEGPU): + print("hostwrite5") + a = FLAMEGPU.environment.getMacroPropertyUInt("a"); + a.set(5); + b = FLAMEGPU.environment.getMacroPropertyUInt("b"); + b[0] = 21; + +class Host_Read_5(pyflamegpu.HostFunctionCallback): + def run(self, FLAMEGPU): + print("hostread5") + result = FLAMEGPU.environment.getMacroPropertyUInt("a"); + assert result == 5; + result2 = FLAMEGPU.environment.getMacroPropertyUInt("b"); + assert result2 == 21; + +Agent_Write_5 = """ +FLAMEGPU_AGENT_FUNCTION(Agent_Write_5, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("a").exchange(5u); + FLAMEGPU->environment.getMacroProperty("b").exchange(21u); + return flamegpu::ALIVE; +} +""" +Agent_Read_Write_5 = """ +FLAMEGPU_AGENT_FUNCTION(Agent_Read_Write_5, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("a", FLAMEGPU->environment.getMacroProperty("a").exchange(0u)); + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("b").exchange(0u)); + return flamegpu::ALIVE; +} +""" +class Host_Agent_Read_5(pyflamegpu.HostFunctionCallback): + def run(self, FLAMEGPU): + agt = FLAMEGPU.agent("test"); + pop = agt.getPopulationData(); + result = pop[0].getVariableUInt("a"); + assert result == 5; + result2 = pop[0].getVariableUInt("b"); + assert result2 == 21; + +class SubCUDAMacroEnvironmentTest(TestCase): + def test_SubWriteHostMasterRead(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exit_always = ExitAlways() + m2.addExitConditionCallback(exit_always); + m2.Environment().newMacroPropertyUInt("a"); + m2.Environment().newMacroPropertyUInt("b"); + hw5 = Host_Write_5(); + m2.addStepFunctionCallback(hw5); + m2.newAgent("test"); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyUInt("a"); + m.Environment().newMacroPropertyUInt("b"); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + senv.autoMap(); + m.newLayer().addSubModel(sm); + hr5 = Host_Read_5(); + m.newLayer().addHostFunctionCallback(hr5); + pop = pyflamegpu.AgentVector(m.newAgent("test"), 1); + + sim = pyflamegpu.CUDASimulation(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + sim.simulate(); + + def test_SubWriteAgentMasterRead(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exit_always = ExitAlways() + m2.addExitConditionCallback(exit_always); + m2.Environment().newMacroPropertyUInt("a"); + m2.Environment().newMacroPropertyUInt("b"); + af = m2.newAgent("test").newRTCFunction("t", Agent_Write_5); + m2.newLayer().addAgentFunction(af); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyUInt("a"); + m.Environment().newMacroPropertyUInt("b"); + agt = m.newAgent("test"); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + senv.autoMapMacroProperties(); + sm.bindAgent("test", "test"); + m.newLayer().addSubModel(sm); + hr5 = Host_Read_5(); + m.newLayer().addHostFunctionCallback(hr5); + pop = pyflamegpu.AgentVector(agt, 1); + + sim = pyflamegpu.CUDASimulation(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + sim.simulate(); + + def test_MasterWriteSubReadHost(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exit_always = ExitAlways() + m2.addExitConditionCallback(exit_always); + m2.Environment().newMacroPropertyUInt("a"); + m2.Environment().newMacroPropertyUInt("b"); + hr5 = Host_Read_5(); + m2.addStepFunctionCallback(hr5); + m2.newAgent("test"); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyUInt("a"); + m.Environment().newMacroPropertyUInt("b"); + agt = m.newAgent("test"); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + senv.autoMap(); + hw5 = Host_Write_5(); + m.newLayer().addHostFunctionCallback(hw5); + m.newLayer().addSubModel(sm); + pop = pyflamegpu.AgentVector(agt, 1); + + sim = pyflamegpu.CUDASimulation(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + sim.simulate(); + + def test_MasterWriteSubReadAgent(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exit_always = ExitAlways() + m2.addExitConditionCallback(exit_always); + m2.Environment().newMacroPropertyUInt("a"); + m2.Environment().newMacroPropertyUInt("b"); + agt = m2.newAgent("test"); + agt.newVariableUInt("a"); + agt.newVariableUInt("b"); + af = agt.newRTCFunction("arw", Agent_Read_Write_5); + m2.newLayer().addAgentFunction(af); + har5 = Host_Agent_Read_5(); + m2.addStepFunctionCallback(har5); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyUInt("a"); + m.Environment().newMacroPropertyUInt("b"); + agt = m.newAgent("test"); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + senv.autoMap(); + sm.bindAgent("test", "test"); + hw5 = Host_Write_5(); + m.newLayer().addHostFunctionCallback(hw5); + m.newLayer().addSubModel(sm); + pop = pyflamegpu.AgentVector(agt, 1); + + sim = pyflamegpu.CUDASimulation(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + sim.simulate(); diff --git a/tests/swig/python/model/test_subenvironment.py b/tests/swig/python/model/test_subenvironment.py index 52808d89e..e71f8debe 100644 --- a/tests/swig/python/model/test_subenvironment.py +++ b/tests/swig/python/model/test_subenvironment.py @@ -119,4 +119,108 @@ def test_AlreadyBound(self): senv.mapProperty("a_", "b"); with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception senv.mapProperty("a2_", "b2"); + + def test_Macro_InvalidNames(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exitcdn = ExitAlways() + m2.addExitConditionCallback(exitcdn); + m2.Environment().newMacroPropertyFloat("a"); + m2.Environment().newMacroPropertyFloat("a2", 2, 3, 4, 5); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyFloat("b"); + m.Environment().newMacroPropertyFloat("b2", 2, 3, 4, 5); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("c", "b"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("c", "b2"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "c"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a2", "c"); + senv.mapMacroProperty("a2", "b2"); + senv.mapMacroProperty("a", "b"); + + def test_Macro_TypesDoNotMatch(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exitcdn = ExitAlways() + m2.addExitConditionCallback(exitcdn); + m2.Environment().newMacroPropertyFloat("a"); + m2.Environment().newMacroPropertyInt("a2"); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyUInt("b"); + m.Environment().newMacroPropertyFloat("b2", 2, 3, 4); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "b"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a2", "b2"); + + def test_Macro_DimensionsDoNotMatch(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exitcdn = ExitAlways() + m2.addExitConditionCallback(exitcdn); + m2.Environment().newMacroPropertyFloat("a", 4, 3, 2, 1); + m2.Environment().newMacroPropertyFloat("a2", 1, 2, 3, 4); + m2.Environment().newMacroPropertyFloat("a3", 1, 2, 3); + m2.Environment().newMacroPropertyFloat("a4", 2, 3, 4); + m2.Environment().newMacroPropertyFloat("a5"); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyFloat("b", 4, 3, 2, 1); + m.Environment().newMacroPropertyFloat("b2", 1, 2, 3, 4); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "b2"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "b3"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "b4"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "b5"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a2", "b"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a3", "b"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a4", "b"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a5", "b"); + senv.mapMacroProperty("a2", "b2"); + senv.mapMacroProperty("a", "b"); + def test_Macro_AlreadyBound(self): + m2 = pyflamegpu.ModelDescription("sub"); + # Define SubModel + exitcdn = ExitAlways() + m2.addExitConditionCallback(exitcdn); + m2.Environment().newMacroPropertyFloat("a"); + m2.Environment().newMacroPropertyFloat("a2", 2); + m2.Environment().newMacroPropertyFloat("a_"); + m2.Environment().newMacroPropertyFloat("a2_", 2); + m = pyflamegpu.ModelDescription("host"); + # Define Model + m.Environment().newMacroPropertyFloat("b"); + m.Environment().newMacroPropertyFloat("b2", 2); + m.Environment().newMacroPropertyFloat("b_"); + m.Environment().newMacroPropertyFloat("b2_", 2); + sm = m.newSubModel("sub", m2); + senv = sm.SubEnvironment(); + senv.mapMacroProperty("a", "b"); + senv.mapMacroProperty("a2", "b2"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a", "b_"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a2", "b2_"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a_", "b"); + with pytest.raises(pyflamegpu.FLAMEGPURuntimeException) as e: # exception::InvalidEnvProperty exception + senv.mapMacroProperty("a2_", "b2"); diff --git a/tests/swig/python/runtime/test_device_macro_property.py b/tests/swig/python/runtime/test_device_macro_property.py new file mode 100644 index 000000000..dc439ec29 --- /dev/null +++ b/tests/swig/python/runtime/test_device_macro_property.py @@ -0,0 +1,760 @@ +""" + Tests of class: DeviceMacroProperty + WriteRead: Check that SEATBELTS catches a read after write in same agent fn + ReadWrite: Check that SEATBELTS catches a write after read in same agent fn + add: Use DeviceAPI operator+=, then read the value back in a subsequent agent function + add2: Use DeviceAPI operator+, and read the returned result + sub: Use DeviceAPI operator-=, then read the value back in a subsequent agent function + sub2: Use DeviceAPI operator-, and read the returned result + preincrement: Use DeviceAPI operator++ (pre), check the results, then read the value back in a subsequent agent function + predecrement: Use DeviceAPI operator-- (pre), check the results, then read the value back in a subsequent agent function + postincrement: Use DeviceAPI operator++ (post), check the results, then read the value back in a subsequent agent function + postdecrement: Use DeviceAPI operator-- (post), check the results, then read the value back in a subsequent agent function + min: Use DeviceAPI min with a value that succeeds, check the results, then read the value back in a subsequent agent function + min2: Use DeviceAPI min with a value that fails, check the results, then read the value back in a subsequent agent function + max: Use DeviceAPI max with a value that succeeds, check the results, then read the value back in a subsequent agent function + max2: Use DeviceAPI max with a value that fails, check the results, then read the value back in a subsequent agent function + cas: Use DeviceAPI cas with a value that succeeds, check the results, then read the value back in a subsequent agent function + cas2: Use DeviceAPI cas with a value that fails, check the results, then read the value back in a subsequent agent function + exchange_int64: Atomic exchange int64, then read the value back in a subsequent agent function + exchange_uint64: Atomic exchange uint64, then read the value back in a subsequent agent function + exchange_int32: Atomic exchange int32, then read the value back in a subsequent agent function + exchange_uint32: Atomic exchange uint32, then read the value back in a subsequent agent function + exchange_double: Atomic exchange double, then read the value back in a subsequent agent function + exchange_float: Atomic exchange float, then read the value back in a subsequent agent function +""" +import pytest +from unittest import TestCase +from pyflamegpu import * + +WriteRead = """ +FLAMEGPU_AGENT_FUNCTION(WriteRead, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +ReadWrite = """ +FLAMEGPU_AGENT_FUNCTION(ReadWrite, flamegpu::MessageNone, flamegpu::MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int"); + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + FLAMEGPU->setVariable("b", t); + return flamegpu::ALIVE; +} +""" +Init_add = """ +FLAMEGPU_AGENT_FUNCTION(Init_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(1); + return flamegpu::ALIVE; +} +""" +Write_add = """ +FLAMEGPU_AGENT_FUNCTION(Write_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + return flamegpu::ALIVE; +} +""" +Read_add = """ +FLAMEGPU_AGENT_FUNCTION(Read_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Write_add2 = """ +FLAMEGPU_AGENT_FUNCTION(Write_add2, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int") + FLAMEGPU->getVariable("a")); + return flamegpu::ALIVE; +} +""" +Init_sub = """ +FLAMEGPU_AGENT_FUNCTION(Init_sub, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(25); + return flamegpu::ALIVE; +} +""" +Write_sub = """ +FLAMEGPU_AGENT_FUNCTION(Write_sub, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") -= FLAMEGPU->getVariable("a"); + return flamegpu::ALIVE; +} +""" +Read_sub = """ +FLAMEGPU_AGENT_FUNCTION(Read_sub, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Write_postincrement = """ +FLAMEGPU_AGENT_FUNCTION(Write_postincrement, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")++); + return flamegpu::ALIVE; +} +""" +Read_increment = """ +FLAMEGPU_AGENT_FUNCTION(Read_increment, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("c", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Write_sub2 = """ +FLAMEGPU_AGENT_FUNCTION(Write_sub2, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int") - FLAMEGPU->getVariable("a")); + return flamegpu::ALIVE; +} +""" +Write_preincrement = """ +FLAMEGPU_AGENT_FUNCTION(Write_preincrement, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", ++FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Write_postdecrement = """ +FLAMEGPU_AGENT_FUNCTION(Write_postdecrement, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")--); + return flamegpu::ALIVE; +} +""" +Write_predecrement = """ +FLAMEGPU_AGENT_FUNCTION(Write_predecrement, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", --FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Write_min = """ +FLAMEGPU_AGENT_FUNCTION(Write_min, flamegpu::MessageNone, flamegpu::MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int").min(FLAMEGPU->getVariable("a")); + FLAMEGPU->setVariable("b", t); + return flamegpu::ALIVE; +} +""" +Write_max = """ +FLAMEGPU_AGENT_FUNCTION(Write_max, flamegpu::MessageNone, flamegpu::MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int").max(FLAMEGPU->getVariable("a")); + FLAMEGPU->setVariable("b", t); + return flamegpu::ALIVE; +} +""" +Write_cas = """ +FLAMEGPU_AGENT_FUNCTION(Write_cas, flamegpu::MessageNone, flamegpu::MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int").CAS(FLAMEGPU->getVariable("a"), 100); + FLAMEGPU->setVariable("b", t); + return flamegpu::ALIVE; +} +""" +Init_int64 = """ +FLAMEGPU_AGENT_FUNCTION(Init_int64, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return flamegpu::ALIVE; +} +""" +Read_int64 = """ +FLAMEGPU_AGENT_FUNCTION(Read_int64, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Init_uint64 = """ +FLAMEGPU_AGENT_FUNCTION(Init_uint64, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15u); + return flamegpu::ALIVE; +} +""" +Read_uint64 = """ +FLAMEGPU_AGENT_FUNCTION(Read_uint64, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Init_int32 = """ +FLAMEGPU_AGENT_FUNCTION(Init_int32, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return flamegpu::ALIVE; +} +""" +Read_int32 = """ +FLAMEGPU_AGENT_FUNCTION(Read_int32, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Init_uint32 = """ +FLAMEGPU_AGENT_FUNCTION(Init_uint32, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15u); + return flamegpu::ALIVE; +} +""" +Read_uint32 = """ +FLAMEGPU_AGENT_FUNCTION(Read_uint32, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Init_double = """ +FLAMEGPU_AGENT_FUNCTION(Init_double, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return flamegpu::ALIVE; +} +""" +Read_double = """ +FLAMEGPU_AGENT_FUNCTION(Read_double, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" +Init_float = """ +FLAMEGPU_AGENT_FUNCTION(Init_float, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return flamegpu::ALIVE; +} +""" +Read_float = """ +FLAMEGPU_AGENT_FUNCTION(Read_float, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +""" + +class DeviceMacroPropertyTest(TestCase): + def test_add(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + initFn = agent.newRTCFunction("init", Init_add); + writeFn = agent.newRTCFunction("write", Write_add); + readFn = agent.newRTCFunction("read", Read_add); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 13 == t_out; + + def test_add2(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + initFn = agent.newRTCFunction("init", Init_add); + writeFn = agent.newRTCFunction("write", Write_add2); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 13 == t_out; + + def test_sub(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_sub); + readFn = agent.newRTCFunction("read", Read_sub); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 13 == t_out; + + def test_sub2(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_sub2); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 13 == t_out; + + def test_postincrement(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_add); + writeFn = agent.newRTCFunction("write", Write_postincrement); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 1 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 2 == t_out2; + + def test_preincrement(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_add); + writeFn = agent.newRTCFunction("write", Write_preincrement); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 2 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 2 == t_out2; + + def test_postdecrement(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_postdecrement); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 25 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 24 == t_out2; + + def test_predecrement(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_predecrement); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 24 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 24 == t_out2; + + def test_predecrement(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_min); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 12 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 12 == t_out2; + + def test_min2(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_min); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 45); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 25 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 25 == t_out2; + + def test_max(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_max); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 45); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 45 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 45 == t_out2; + + def test_max2(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_max); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 25 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 25 == t_out2; + + def test_cas(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_cas); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 25); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 25 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 100 == t_out2; + + def test_cas2(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + agent.newVariableUInt("c"); + initFn = agent.newRTCFunction("init", Init_sub); + writeFn = agent.newRTCFunction("write", Write_cas); + readFn = agent.newRTCFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt("b"); + assert 25 == t_out; + t_out2 = population.at(0).getVariableUInt("c"); + assert 25 == t_out2; + + def test_exchange_int64(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyInt64("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableInt64("b"); + initFn = agent.newRTCFunction("init", Init_int64); + readFn = agent.newRTCFunction("read", Read_int64); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableInt64("b"); + assert 15 == t_out; + + def test_exchange_uint64(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt64("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt64("b"); + initFn = agent.newRTCFunction("init", Init_uint64); + readFn = agent.newRTCFunction("read", Read_uint64); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt64("b"); + assert 15 == t_out; + + def test_exchange_int32(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyInt32("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableInt32("b"); + initFn = agent.newRTCFunction("init", Init_int32); + readFn = agent.newRTCFunction("read", Read_int32); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableInt32("b"); + assert 15 == t_out; + + def test_exchange_uint32(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt32("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt32("b"); + initFn = agent.newRTCFunction("init", Init_uint32); + readFn = agent.newRTCFunction("read", Read_uint32); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableUInt32("b"); + assert 15 == t_out; + + def test_exchange_double(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyDouble("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableDouble("b"); + initFn = agent.newRTCFunction("init", Init_double); + readFn = agent.newRTCFunction("read", Read_double); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableDouble("b"); + assert 15.0 == t_out; + + def test_exchange_float(self): + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyFloat("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableFloat("b"); + initFn = agent.newRTCFunction("init", Init_float); + readFn = agent.newRTCFunction("read", Read_float); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + population = pyflamegpu.AgentVector(agent, 1); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + t_out = population.at(0).getVariableFloat("b"); + assert 15.0 == t_out; + + def test_WriteRead(self): + if not pyflamegpu.SEATBELTS: + pytest.skip("Test requires SEATBELTS to be ON") + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + WriteReadFn = agent.newRTCFunction("WriteRead", WriteRead); + model.newLayer().addAgentFunction(WriteReadFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + try: + cudaSimulation.simulate(); # DeviceError + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + def test_ReadWrite(self): + if not pyflamegpu.SEATBELTS: + pytest.skip("Test requires SEATBELTS to be ON") + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("a"); + agent.newVariableUInt("b"); + ReadWriteFn = agent.newRTCFunction("ReadWrite", ReadWrite); + model.newLayer().addAgentFunction(ReadWriteFn); + population = pyflamegpu.AgentVector(agent, 1); + population[0].setVariableUInt("a", 12); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + try: + cudaSimulation.simulate(); # DeviceError + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; diff --git a/tests/swig/python/runtime/test_host_macro_property.py b/tests/swig/python/runtime/test_host_macro_property.py new file mode 100644 index 000000000..00a36aec9 --- /dev/null +++ b/tests/swig/python/runtime/test_host_macro_property.py @@ -0,0 +1,578 @@ +import pytest +from unittest import TestCase +from pyflamegpu import * + +""" + Tests of class: HostMacroProperty + ReadTest: Test that HostMacroProperty can read data (written via agent fn's) + WriteTest: Test that HostMacroProperty can write data (read via agent fn's) + ZeroTest: Test that HostMacroProperty can zero data (read via agent fn's) +""" + +TEST_DIMS = [2, 3, 4, 5]; + +class HostRead(pyflamegpu.HostFunctionCallback): + def run(self,FLAMEGPU): + t = FLAMEGPU.environment.getMacroPropertyUInt("int"); + a = 0; + for i in t: + for j in i: + for k in j: + for w in k: + assert w == a + a+=1; + # Test both iteration styles + a = 0; + for i in range(TEST_DIMS[0]): + for j in range(TEST_DIMS[1]): + for k in range(TEST_DIMS[2]): + for w in range(TEST_DIMS[3]): + assert t[i][j][k][w] == a + a+=1; + +AgentWrite = """ +FLAMEGPU_AGENT_FUNCTION(AgentWrite, flamegpu::MessageNone, flamegpu::MessageNone) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + const unsigned int i = FLAMEGPU->getVariable("i"); + const unsigned int j = FLAMEGPU->getVariable("j"); + const unsigned int k = FLAMEGPU->getVariable("k"); + const unsigned int w = FLAMEGPU->getVariable("w"); + t[i][j][k][w].exchange(FLAMEGPU->getVariable("a")); + return flamegpu::ALIVE; +} +""" + +class HostWrite(pyflamegpu.HostFunctionCallback): + def run(self,FLAMEGPU): + t = FLAMEGPU.environment.getMacroPropertyUInt("int"); + a = 0; + for i in t: + for j in i: + for k in j: + for w in k: + w.set(a); + a+=1; + +AgentRead = """ +FLAMEGPU_AGENT_FUNCTION(AgentRead, flamegpu::MessageNone, flamegpu::MessageNone) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + const unsigned int i = FLAMEGPU->getVariable("i"); + const unsigned int j = FLAMEGPU->getVariable("j"); + const unsigned int k = FLAMEGPU->getVariable("k"); + const unsigned int w = FLAMEGPU->getVariable("w"); + if (t[i][j][k][w] == FLAMEGPU->getVariable("a")) { + FLAMEGPU->setVariable("a", 1); + } else { + FLAMEGPU->setVariable("a", 0); + } + return flamegpu::ALIVE; +} +""" + +class HostZero(pyflamegpu.HostFunctionCallback): + def run(self,FLAMEGPU): + t = FLAMEGPU.environment.getMacroPropertyUInt("int"); + t.zero(); + +AgentReadZero = """ +FLAMEGPU_AGENT_FUNCTION(AgentReadZero, flamegpu::MessageNone, flamegpu::MessageNone) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + const unsigned int i = FLAMEGPU->getVariable("i"); + const unsigned int j = FLAMEGPU->getVariable("j"); + const unsigned int k = FLAMEGPU->getVariable("k"); + const unsigned int w = FLAMEGPU->getVariable("w"); + if (t[i][j][k][w] == 0) { + FLAMEGPU->setVariable("a", 1); + } else { + FLAMEGPU->setVariable("a", 0); + } + return flamegpu::ALIVE; +} +""" + +class HostArithmeticInit(pyflamegpu.HostFunctionCallback): + def run(self,FLAMEGPU): + FLAMEGPU.environment.getMacroPropertyInt("int")[0] = 10; + FLAMEGPU.environment.getMacroPropertyUInt("uint")[0] = 10; + FLAMEGPU.environment.getMacroPropertyInt8("int8")[0] = 10; + FLAMEGPU.environment.getMacroPropertyUInt8("uint8")[0] = 10; + FLAMEGPU.environment.getMacroPropertyInt64("int64")[0] = 10; + FLAMEGPU.environment.getMacroPropertyUInt64("uint64")[0] = 10; + FLAMEGPU.environment.getMacroPropertyID("id")[0] = 10; + FLAMEGPU.environment.getMacroPropertyFloat("float")[0] = 10; + FLAMEGPU.environment.getMacroPropertyDouble("double").set(10); # alt + +class HostArithmetic(pyflamegpu.HostFunctionCallback): + def run(self,FLAMEGPU): + # int + t = FLAMEGPU.environment.getMacroPropertyInt("int"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # uint + t = FLAMEGPU.environment.getMacroPropertyUInt("uint"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # int8 + t = FLAMEGPU.environment.getMacroPropertyInt8("int8"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # uint8 + t = FLAMEGPU.environment.getMacroPropertyUInt8("uint8"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # int64 + t = FLAMEGPU.environment.getMacroPropertyInt64("int64"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # uint64 + t = FLAMEGPU.environment.getMacroPropertyUInt64("uint64"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # id + t = FLAMEGPU.environment.getMacroPropertyID("id"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10 / 5; + assert t / 3 == 10 / 3; + assert t / 3.0 == 10 / 3.0; + assert t + 5 == 10 + 5; + assert t + 3.0 == 10 + 3.0; + assert t - 3 == 10 - 3; + assert t - 3.0 == 10 - 3.0; + assert t * 5 == 10 * 5; + assert t * 3.0 == 10 * 3.0; + assert t % 5 == 10 % 5; + assert t % 3 == 10 % 3; + assert t // 5 == 10 // 5; + assert t // 3 == 10 // 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t %= 4; + assert t == 2; + try: + t /= 2; # int does no support true div, as would convert to float + except pyflamegpu.FLAMEGPURuntimeException: + assert True; + else: + assert False; + + # float + t = FLAMEGPU.environment.getMacroPropertyFloat("float"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10.0 / 5; + assert t / 3 == 10.0 / 3; + assert t + 5 == 10.0 + 5; + assert t - 3 == 10.0 - 3; + assert t * 5 == 10.0 * 5; + assert t // 5 == 10.0 // 5; + assert t // 3 == 10.0 // 3; + assert t % 5 == 10.0 % 5; + assert t % 3 == 10.0 % 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t /= 2; + assert t == 5; + t.set(10); + t %= 4.0; + assert t == 2.0; + + # double + t = FLAMEGPU.environment.getMacroPropertyDouble("double"); + # assert t++ == 10; # Python does not support overloading increment operator + # assert ++t == 12; # Python does not support overloading increment operator + # assert t-- == 12; # Python does not support overloading increment operator + # assert --t == 10; # Python does not support overloading increment operator + assert t / 5 == 10.0 / 5; + assert t / 3 == 10.0 / 3; + assert t + 5 == 10.0 + 5; + assert t - 3 == 10.0 - 3; + assert t * 5 == 10.0 * 5; + assert t // 5 == 10.0 // 5; + assert t // 3 == 10.0 // 3; + assert t % 5 == 10.0 % 5; + assert t % 3 == 10.0 % 3; + assert t == 10; + t += 10; + assert t == 20; + t -= 10; + assert t == 10; + t *= 2; + assert t == 20; + t //= 2; + assert t == 10; + t /= 2; + assert t == 5; + t.set(10); + t %= 4.0; + assert t == 2.0; + +class HostMacroPropertyTest(TestCase): + def test_ReadTest(self): + # Fill MacroProperty with DeviceAPI + # Test values match expected value with HostAPI and test in-place + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int", 2, 3, 4, 5); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("i"); + agent.newVariableUInt("j"); + agent.newVariableUInt("k"); + agent.newVariableUInt("w"); + agent.newVariableUInt("a"); + t = agent.newRTCFunction("agentwrite", AgentWrite); + model.newLayer().addAgentFunction(t); + model.newLayer().addHostFunctionCallback(HostRead().__disown__()); + total_agents = TEST_DIMS[0] * TEST_DIMS[1] * TEST_DIMS[2] * TEST_DIMS[3]; + population = pyflamegpu.AgentVector(agent, total_agents); + a = 0; + for i in range (TEST_DIMS[0]): + for j in range (TEST_DIMS[1]): + for k in range (TEST_DIMS[2]): + for w in range (TEST_DIMS[3]): + p = population[a]; + p.setVariableUInt("a", a); + p.setVariableUInt("i", i); + p.setVariableUInt("j", j); + p.setVariableUInt("k", k); + p.setVariableUInt("w", w); + a += 1; + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + + def test_WriteTest(self): + # Fill MacroProperty with HostAPI + # Test values match expected value with DeviceAPI + # Write results back to agent variable, and check at the end + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int", 2, 3, 4, 5); + model.Environment().newMacroPropertyUInt("plusequal"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("i"); + agent.newVariableUInt("j"); + agent.newVariableUInt("k"); + agent.newVariableUInt("w"); + agent.newVariableUInt("a"); + t = agent.newRTCFunction("agentread", AgentRead); + model.newLayer().addHostFunctionCallback(HostWrite().__disown__()); + model.newLayer().addAgentFunction(t); + total_agents = TEST_DIMS[0] * TEST_DIMS[1] * TEST_DIMS[2] * TEST_DIMS[3]; + population = pyflamegpu.AgentVector(agent, total_agents); + a = 0; + for i in range (TEST_DIMS[0]): + for j in range (TEST_DIMS[1]): + for k in range (TEST_DIMS[2]): + for w in range (TEST_DIMS[3]): + p = population[a]; + p.setVariableUInt("a", a); + p.setVariableUInt("i", i); + p.setVariableUInt("j", j); + p.setVariableUInt("k", k); + p.setVariableUInt("w", w); + a += 1; + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + # Check results + correct = 0; + for p in population: + correct += 1 if p.getVariableUInt("a") == 1 else 0; + + assert correct == total_agents; + + def test_ZeroTest(self): + # Fill MacroProperty with HostAPI + # Test values match expected value with DeviceAPI + # Write results back to agent variable, and check at the end + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyUInt("int", 2, 3, 4, 5); + model.Environment().newMacroPropertyUInt("plusequal"); + # Setup agent fn + agent = model.newAgent("agent"); + agent.newVariableUInt("i"); + agent.newVariableUInt("j"); + agent.newVariableUInt("k"); + agent.newVariableUInt("w"); + agent.newVariableUInt("a"); + t1 = agent.newRTCFunction("agentwrite", AgentWrite); + t2 = agent.newRTCFunction("agentread", AgentReadZero); + model.newLayer().addAgentFunction(t1); + model.newLayer().addHostFunctionCallback(HostZero().__disown__()); + model.newLayer().addAgentFunction(t2); + total_agents = TEST_DIMS[0] * TEST_DIMS[1] * TEST_DIMS[2] * TEST_DIMS[3]; + population = pyflamegpu.AgentVector(agent, total_agents); + a = 0; + for i in range (TEST_DIMS[0]): + for j in range (TEST_DIMS[1]): + for k in range (TEST_DIMS[2]): + for w in range (TEST_DIMS[3]): + p = population[a]; + p.setVariableUInt("a", a); + p.setVariableUInt("i", i); + p.setVariableUInt("j", j); + p.setVariableUInt("k", k); + p.setVariableUInt("w", w); + a += 1; + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + cudaSimulation.simulate(); + cudaSimulation.getPopulationData(population); + # Check results + correct = 0; + for p in population: + correct += 1 if p.getVariableUInt("a") == 1 else 0; + + assert correct == total_agents; + + def test_ArithmeticTest(self): + # Create single macro property for each type + # Fill MacroProperties with HostAPI to a known value + # Use all airthmetic ops, and test values match expected value with DeviceAPI + model = pyflamegpu.ModelDescription("device_env_test"); + # Setup environment + model.Environment().newMacroPropertyInt("int"); + model.Environment().newMacroPropertyUInt("uint"); + model.Environment().newMacroPropertyInt8("int8"); + model.Environment().newMacroPropertyUInt8("uint8"); + model.Environment().newMacroPropertyInt64("int64"); + model.Environment().newMacroPropertyUInt64("uint64"); + model.Environment().newMacroPropertyID("id"); + model.Environment().newMacroPropertyFloat("float"); + model.Environment().newMacroPropertyDouble("double"); + # Setup agent fn + model.newAgent("agent"); + model.newLayer().addHostFunctionCallback(HostArithmeticInit().__disown__()); + model.newLayer().addHostFunctionCallback(HostArithmetic().__disown__()); + # Do Sim + cudaSimulation = pyflamegpu.CUDASimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.simulate(); + \ No newline at end of file diff --git a/tests/test_cases/gpu/test_cuda_submacroenvironment.cu b/tests/test_cases/gpu/test_cuda_submacroenvironment.cu new file mode 100644 index 000000000..880938dc5 --- /dev/null +++ b/tests/test_cases/gpu/test_cuda_submacroenvironment.cu @@ -0,0 +1,169 @@ +/** + * Tests of submodel functionality of class: CUDAMacroEnvironment + * ReadTest: Test that HostMacroProperty can read data (written via agent fn's) + * WriteTest: Test that HostMacroProperty can write data (read via agent fn's) + * ZeroTest: Test that HostMacroProperty can zero data (read via agent fn's) + */ + +#include + +#include "flamegpu/flamegpu.h" + +#include "gtest/gtest.h" + +namespace flamegpu { +namespace { +FLAMEGPU_EXIT_CONDITION(ExitAlways) { + return EXIT; +} +FLAMEGPU_STEP_FUNCTION(Host_Write_5) { + FLAMEGPU->environment.getMacroProperty("a") = 5; + FLAMEGPU->environment.getMacroProperty("b") = 21; +} +FLAMEGPU_HOST_FUNCTION(Host_Read_5) { + const unsigned int result = FLAMEGPU->environment.getMacroProperty("a"); + EXPECT_EQ(result, 5u); + const unsigned int result2 = FLAMEGPU->environment.getMacroProperty("b"); + EXPECT_EQ(result2, 21u); +} +FLAMEGPU_AGENT_FUNCTION(Agent_Write_5, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("a").exchange(5u); + FLAMEGPU->environment.getMacroProperty("b").exchange(21u); + return flamegpu::ALIVE; +} +TEST(SubCUDAMacroEnvironmentTest, SubWriteHostMasterRead) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("b"); + m2.addStepFunction(Host_Write_5); + m2.newAgent("test"); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("a"); + m.Environment().newMacroProperty("b"); + } + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + senv.autoMap(); + m.newLayer().addSubModel(sm); + m.newLayer().addHostFunction(Host_Read_5); + AgentVector pop(m.newAgent("test"), 1); + + CUDASimulation sim(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + EXPECT_NO_THROW(sim.simulate()); +} +TEST(SubCUDAMacroEnvironmentTest, SubWriteAgentMasterRead) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("b"); + m2.newAgent("test").newFunction("t", Agent_Write_5); + m2.newLayer().addAgentFunction(Agent_Write_5); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("a"); + m.Environment().newMacroProperty("b"); + } + auto &agt = m.newAgent("test"); + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + senv.autoMapMacroProperties(); + sm.bindAgent("test", "test"); + m.newLayer().addSubModel(sm); + m.newLayer().addHostFunction(Host_Read_5); + AgentVector pop(agt, 1); + + CUDASimulation sim(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + EXPECT_NO_THROW(sim.simulate()); +} +TEST(SubCUDAMacroEnvironmentTest, MasterWriteSubReadHost) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("b"); + m2.addStepFunction(Host_Read_5); + m2.newAgent("test"); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("a"); + m.Environment().newMacroProperty("b"); + } + auto& agt = m.newAgent("test"); + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + senv.autoMap(); + m.newLayer().addHostFunction(Host_Write_5); + m.newLayer().addSubModel(sm); + AgentVector pop(agt, 1); + + CUDASimulation sim(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + EXPECT_NO_THROW(sim.simulate()); +} +FLAMEGPU_AGENT_FUNCTION(Agent_Read_Write_5, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("a", FLAMEGPU->environment.getMacroProperty("a").exchange(0u)); + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("b").exchange(0u)); + return flamegpu::ALIVE; +} +FLAMEGPU_HOST_FUNCTION(Host_Agent_Read_5) { + auto agt = FLAMEGPU->agent("test"); + DeviceAgentVector pop = agt.getPopulationData(); + const unsigned int result = pop[0].getVariable("a"); + EXPECT_EQ(result, 5u); + const unsigned int result2 = pop[0].getVariable("b"); + EXPECT_EQ(result2, 21u); +} +TEST(SubCUDAMacroEnvironmentTest, MasterWriteSubReadAgent) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("b"); + auto &agt = m2.newAgent("test"); + agt.newVariable("a"); + agt.newVariable("b"); + agt.newFunction("arw", Agent_Read_Write_5); + m2.newLayer().addAgentFunction(Agent_Read_Write_5); + m2.addStepFunction(Host_Agent_Read_5); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("a"); + m.Environment().newMacroProperty("b"); + } + auto& agt = m.newAgent("test"); + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + senv.autoMap(); + sm.bindAgent("test", "test"); + m.newLayer().addHostFunction(Host_Write_5); + m.newLayer().addSubModel(sm); + AgentVector pop(agt, 1); + + CUDASimulation sim(m); + sim.SimulationConfig().steps = 1; + sim.setPopulationData(pop); + EXPECT_NO_THROW(sim.simulate()); +} +} // namespace +} // namespace flamegpu diff --git a/tests/test_cases/model/test_subenvironment.cu b/tests/test_cases/model/test_subenvironment.cu index 4b2fd34e2..188133933 100644 --- a/tests/test_cases/model/test_subenvironment.cu +++ b/tests/test_cases/model/test_subenvironment.cu @@ -120,5 +120,105 @@ TEST(SubEnvironmentDescriptionTest, AlreadyBound) { EXPECT_THROW(senv.mapProperty("a_", "b"), exception::InvalidEnvProperty); EXPECT_THROW(senv.mapProperty("a2_", "b2"), exception::InvalidEnvProperty); } +TEST(SubEnvironmentDescriptionTest, Macro_InvalidNames) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("a2"); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("b"); + m.Environment().newMacroProperty("b2"); + } + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + EXPECT_THROW(senv.mapMacroProperty("c", "b"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("c", "b2"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a", "c"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a2", "c"), exception::InvalidEnvProperty); + EXPECT_NO_THROW(senv.mapMacroProperty("a2", "b2")); + EXPECT_NO_THROW(senv.mapMacroProperty("a", "b")); +} +TEST(SubEnvironmentDescriptionTest, Macro_TypesDoNotMatch) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("a2"); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("b"); + m.Environment().newMacroProperty("b2"); + } + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + EXPECT_THROW(senv.mapMacroProperty("a", "b"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a2", "b2"), exception::InvalidEnvProperty); +} +TEST(SubEnvironmentDescriptionTest, Macro_DimensionsDoNotMatch) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("a2"); + m2.Environment().newMacroProperty("a3"); + m2.Environment().newMacroProperty("a4"); + m2.Environment().newMacroProperty("a5"); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("b"); + m.Environment().newMacroProperty("b2"); + } + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + EXPECT_THROW(senv.mapMacroProperty("a", "b2"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a", "b3"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a", "b4"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a", "b5"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a2", "b"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a3", "b"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a4", "b"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a5", "b"), exception::InvalidEnvProperty); + EXPECT_NO_THROW(senv.mapMacroProperty("a2", "b2")); + EXPECT_NO_THROW(senv.mapMacroProperty("a", "b")); +} +TEST(SubEnvironmentDescriptionTest, Macro_AlreadyBound) { + ModelDescription m2("sub"); + { + // Define SubModel + m2.addExitCondition(ExitAlways); + m2.Environment().newMacroProperty("a"); + m2.Environment().newMacroProperty("a2"); + m2.Environment().newMacroProperty("a_"); + m2.Environment().newMacroProperty("a2_"); + } + ModelDescription m("host"); + { + // Define Model + m.Environment().newMacroProperty("b"); + m.Environment().newMacroProperty("b2"); + m.Environment().newMacroProperty("b_"); + m.Environment().newMacroProperty("b2_"); + } + // Missing exit condition + auto& sm = m.newSubModel("sub", m2); + auto& senv = sm.SubEnvironment(); + EXPECT_NO_THROW(senv.mapMacroProperty("a", "b")); + EXPECT_NO_THROW(senv.mapMacroProperty("a2", "b2")); + EXPECT_THROW(senv.mapMacroProperty("a", "b_"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a2", "b2_"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a_", "b"), exception::InvalidEnvProperty); + EXPECT_THROW(senv.mapMacroProperty("a2_", "b2"), exception::InvalidEnvProperty); +} }; // namespace test_sub_environment_description } // namespace flamegpu diff --git a/tests/test_cases/runtime/test_agent_environment.cu b/tests/test_cases/runtime/test_device_environment.cu similarity index 95% rename from tests/test_cases/runtime/test_agent_environment.cu rename to tests/test_cases/runtime/test_device_environment.cu index 0c041c06a..4f7da972a 100644 --- a/tests/test_cases/runtime/test_agent_environment.cu +++ b/tests/test_cases/runtime/test_device_environment.cu @@ -185,7 +185,7 @@ class MiniSim { * This defines a common fixture used as a base for all test cases in the file * @see https://github.com/google/googletest/blob/master/googletest/samples/sample5_unittest.cc */ -class AgentEnvironmentTest : public testing::Test { +class DeviceEnvironmentTest : public testing::Test { protected: void SetUp() override { ms = new MiniSim(); @@ -224,7 +224,7 @@ class AgentEnvironmentTest : public testing::Test { } // namespace -TEST_F(AgentEnvironmentTest, Get_float) { +TEST_F(DeviceEnvironmentTest, Get_float) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_float); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -239,7 +239,7 @@ TEST_F(AgentEnvironmentTest, Get_float) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_double) { +TEST_F(DeviceEnvironmentTest, Get_double) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_double); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -254,7 +254,7 @@ TEST_F(AgentEnvironmentTest, Get_double) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_int8_t) { +TEST_F(DeviceEnvironmentTest, Get_int8_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_int8_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -269,7 +269,7 @@ TEST_F(AgentEnvironmentTest, Get_int8_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_uint8_t) { +TEST_F(DeviceEnvironmentTest, Get_uint8_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_uint8_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -284,7 +284,7 @@ TEST_F(AgentEnvironmentTest, Get_uint8_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_int16_t) { +TEST_F(DeviceEnvironmentTest, Get_int16_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_int16_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -299,7 +299,7 @@ TEST_F(AgentEnvironmentTest, Get_int16_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_uint16_t) { +TEST_F(DeviceEnvironmentTest, Get_uint16_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_uint16_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -314,7 +314,7 @@ TEST_F(AgentEnvironmentTest, Get_uint16_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_int32_t) { +TEST_F(DeviceEnvironmentTest, Get_int32_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_int32_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -329,7 +329,7 @@ TEST_F(AgentEnvironmentTest, Get_int32_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_uint32_t) { +TEST_F(DeviceEnvironmentTest, Get_uint32_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_uint32_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -344,7 +344,7 @@ TEST_F(AgentEnvironmentTest, Get_uint32_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_int64_t) { +TEST_F(DeviceEnvironmentTest, Get_int64_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_int64_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -359,7 +359,7 @@ TEST_F(AgentEnvironmentTest, Get_int64_t) { EXPECT_EQ(true, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_uint64_t) { +TEST_F(DeviceEnvironmentTest, Get_uint64_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_uint64_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -375,7 +375,7 @@ TEST_F(AgentEnvironmentTest, Get_uint64_t) { EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_float) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_float) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_float); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -390,7 +390,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_float) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_double) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_double) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_double); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -405,7 +405,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_double) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_int8_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_int8_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_int8_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -420,7 +420,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_int8_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_uint8_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_uint8_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_uint8_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -435,7 +435,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_uint8_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_int16_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_int16_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_int16_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -450,7 +450,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_int16_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_uint16_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_uint16_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_uint16_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -465,7 +465,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_uint16_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_uint32_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_uint32_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_uint32_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -480,7 +480,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_uint32_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_int32_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_int32_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_int32_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -495,7 +495,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_int32_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_uint64_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_uint64_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_uint64_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -510,7 +510,7 @@ TEST_F(AgentEnvironmentTest, Get_arrayElement_uint64_t) { EXPECT_EQ(false, _bool_out); EXPECT_EQ(cudaGetLastError(), CUDA_SUCCESS); } -TEST_F(AgentEnvironmentTest, Get_arrayElement_int64_t) { +TEST_F(DeviceEnvironmentTest, Get_arrayElement_int64_t) { // Setup agent fn AgentFunctionDescription &deviceFn = ms->agent.newFunction("device_function", get_arrayElement_int64_t); LayerDescription &devicefn_layer = ms->model.newLayer("devicefn_layer"); @@ -533,7 +533,7 @@ FLAMEGPU_AGENT_FUNCTION(get_array_glm, MessageNone, MessageNone) { FLAMEGPU->setVariable("k", 2, t[2]); return ALIVE; } -TEST_F(AgentEnvironmentTest, Get_array_glm) { +TEST_F(DeviceEnvironmentTest, Get_array_glm) { // Setup agent fn ms->agent.newVariable("k"); AgentFunctionDescription& deviceFn = ms->agent.newFunction("device_function", get_array_glm); diff --git a/tests/test_cases/runtime/test_device_macro_property.cu b/tests/test_cases/runtime/test_device_macro_property.cu new file mode 100644 index 000000000..486ac0fd4 --- /dev/null +++ b/tests/test_cases/runtime/test_device_macro_property.cu @@ -0,0 +1,777 @@ +/** + * Tests of class: DeviceMacroProperty + * WriteRead: Check that SEATBELTS catches a read after write in same agent fn + * ReadWrite: Check that SEATBELTS catches a write after read in same agent fn + * add: Use DeviceAPI operator+=, then read the value back in a subsequent agent function + * add2: Use DeviceAPI operator+, and read the returned result + * sub: Use DeviceAPI operator-=, then read the value back in a subsequent agent function + * sub2: Use DeviceAPI operator-, and read the returned result + * preincrement: Use DeviceAPI operator++ (pre), check the results, then read the value back in a subsequent agent function + * predecrement: Use DeviceAPI operator-- (pre), check the results, then read the value back in a subsequent agent function + * postincrement: Use DeviceAPI operator++ (post), check the results, then read the value back in a subsequent agent function + * postdecrement: Use DeviceAPI operator-- (post), check the results, then read the value back in a subsequent agent function + * min: Use DeviceAPI min with a value that succeeds, check the results, then read the value back in a subsequent agent function + * min2: Use DeviceAPI min with a value that fails, check the results, then read the value back in a subsequent agent function + * max: Use DeviceAPI max with a value that succeeds, check the results, then read the value back in a subsequent agent function + * max2: Use DeviceAPI max with a value that fails, check the results, then read the value back in a subsequent agent function + * cas: Use DeviceAPI cas with a value that succeeds, check the results, then read the value back in a subsequent agent function + * cas2: Use DeviceAPI cas with a value that fails, check the results, then read the value back in a subsequent agent function + * exchange_int64: Atomic exchange int64, then read the value back in a subsequent agent function + * exchange_uint64: Atomic exchange uint64, then read the value back in a subsequent agent function + * exchange_int32: Atomic exchange int32, then read the value back in a subsequent agent function + * exchange_uint32: Atomic exchange uint32, then read the value back in a subsequent agent function + * exchange_double: Atomic exchange double, then read the value back in a subsequent agent function + * exchange_float: Atomic exchange float, then read the value back in a subsequent agent function + */ +#include "flamegpu/flamegpu.h" + +#include "gtest/gtest.h" + +namespace flamegpu { +namespace { +FLAMEGPU_AGENT_FUNCTION(WriteRead, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +#if !defined(SEATBELTS) || SEATBELTS +TEST(DeviceMacroPropertyTest, WriteRead) { +#else +TEST(DeviceMacroPropertyTest, DISABLED_WriteRead) { +#endif + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& WriteReadFn = agent.newFunction("WriteRead", WriteRead); + model.newLayer().addAgentFunction(WriteReadFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + EXPECT_THROW(cudaSimulation.simulate(), flamegpu::exception::DeviceError); +} +FLAMEGPU_AGENT_FUNCTION(ReadWrite, flamegpu::MessageNone, flamegpu::MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int"); + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + FLAMEGPU->setVariable("b", t); + return flamegpu::ALIVE; +} +#if !defined(SEATBELTS) || SEATBELTS +TEST(DeviceMacroPropertyTest, ReadWrite) { +#else +TEST(DeviceMacroPropertyTest, DISABLED_ReadWrite) { +#endif + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& ReadWriteFn = agent.newFunction("ReadWrite", ReadWrite); + model.newLayer().addAgentFunction(ReadWriteFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + EXPECT_THROW(cudaSimulation.simulate(), flamegpu::exception::DeviceError); +} + +FLAMEGPU_AGENT_FUNCTION(Init_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(1); + return flamegpu::ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Write_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + return flamegpu::ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +TEST(DeviceMacroPropertyTest, add) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_add); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_add); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_add); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(13u, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Write_add2, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int") + FLAMEGPU->getVariable("a")); + return flamegpu::ALIVE; +} +TEST(DeviceMacroPropertyTest, add2) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_add); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_add2); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(13u, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Init_sub, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(25); + return flamegpu::ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Write_sub, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") -= FLAMEGPU->getVariable("a"); + return flamegpu::ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_sub, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +TEST(DeviceMacroPropertyTest, sub) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_sub); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_sub); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(13u, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Write_sub2, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int") - FLAMEGPU->getVariable("a")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, sub2) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_sub2); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(13u, t_out); +} + +FLAMEGPU_AGENT_FUNCTION(Write_postincrement, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")++); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_increment, MessageNone, MessageNone) { + FLAMEGPU->setVariable("c", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, postincrement) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_add); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_postincrement); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(1u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(2u, t_out2); +} +FLAMEGPU_AGENT_FUNCTION(Write_preincrement, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", ++FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, preincrement) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_add); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_preincrement); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(2u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(2u, t_out2); +} +FLAMEGPU_AGENT_FUNCTION(Write_postdecrement, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")--); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, postdecrement) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_postdecrement); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(25u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(24u, t_out2); +} +FLAMEGPU_AGENT_FUNCTION(Write_predecrement, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", --FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, predecrement) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_predecrement); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(24u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(24u, t_out2); +} + +FLAMEGPU_AGENT_FUNCTION(Write_min, MessageNone, MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int").min(FLAMEGPU->getVariable("a")); + FLAMEGPU->setVariable("b", t); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, min) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_min); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(12u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(12u, t_out2); +} +TEST(DeviceMacroPropertyTest, min2) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_min); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 45u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(25u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(25u, t_out2); +} +FLAMEGPU_AGENT_FUNCTION(Write_max, MessageNone, MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int").max(FLAMEGPU->getVariable("a")); + FLAMEGPU->setVariable("b", t); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, max) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_max); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 45u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(45u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(45u, t_out2); +} +TEST(DeviceMacroPropertyTest, max2) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_max); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(25u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(25u, t_out2); +} +FLAMEGPU_AGENT_FUNCTION(Write_cas, MessageNone, MessageNone) { + unsigned int t = FLAMEGPU->environment.getMacroProperty("int").CAS(FLAMEGPU->getVariable("a"), 100); + FLAMEGPU->setVariable("b", t); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, cas) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_cas); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 25u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(25u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(100u, t_out2); +} +TEST(DeviceMacroPropertyTest, cas2) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + agent.newVariable("c"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_sub); + AgentFunctionDescription& writeFn = agent.newFunction("write", Write_cas); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_increment); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(25u, t_out); + const unsigned int t_out2 = population.at(0).getVariable("c"); + ASSERT_EQ(25u, t_out2); +} + +FLAMEGPU_AGENT_FUNCTION(Init_int64, MessageNone, MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_int64, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, exchange_int64) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_int64); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_int64); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const int64_t t_out = population.at(0).getVariable("b"); + ASSERT_EQ(15, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Init_uint64, MessageNone, MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15u); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_uint64, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, exchange_uint64) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_uint64); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_uint64); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const uint64_t t_out = population.at(0).getVariable("b"); + ASSERT_EQ(15u, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Init_int32, MessageNone, MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_int32, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, exchange_int32) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_int32); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_int32); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const int32_t t_out = population.at(0).getVariable("b"); + ASSERT_EQ(15, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Init_uint32, MessageNone, MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15u); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_uint32, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, exchange_uint32) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_uint32); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_uint32); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const uint32_t t_out = population.at(0).getVariable("b"); + ASSERT_EQ(15u, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Init_double, MessageNone, MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_double, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, exchange_double) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_double); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_double); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const double t_out = population.at(0).getVariable("b"); + ASSERT_EQ(15.0, t_out); +} +FLAMEGPU_AGENT_FUNCTION(Init_float, MessageNone, MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(15); + return ALIVE; +} +FLAMEGPU_AGENT_FUNCTION(Read_float, MessageNone, MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return ALIVE; +} +TEST(DeviceMacroPropertyTest, exchange_float) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newFunction("init", Init_float); + AgentFunctionDescription& readFn = agent.newFunction("read", Read_float); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const float t_out = population.at(0).getVariable("b"); + ASSERT_EQ(15.0f, t_out); +} + +const char* WriteRead_func = R"###( +FLAMEGPU_AGENT_FUNCTION(WriteRead, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +)###"; +#if !defined(SEATBELTS) || SEATBELTS +TEST(DeviceMacroPropertyTest, RTC_WriteRead) { +#else +TEST(DeviceMacroPropertyTest, DISABLED_RTC_WriteRead) { +#endif + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& WriteReadFn = agent.newRTCFunction("WriteRead", WriteRead_func); + model.newLayer().addAgentFunction(WriteReadFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + EXPECT_THROW(cudaSimulation.simulate(), flamegpu::exception::DeviceError); +} + +const char* Init_add_func = R"###( +FLAMEGPU_AGENT_FUNCTION(Init_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int").exchange(1); + return flamegpu::ALIVE; +} +)###"; +const char* Write_add_func = R"###( +FLAMEGPU_AGENT_FUNCTION(Write_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->environment.getMacroProperty("int") += FLAMEGPU->getVariable("a"); + return flamegpu::ALIVE; +} +)###"; +const char* Read_add_func = R"###( +FLAMEGPU_AGENT_FUNCTION(Read_add, flamegpu::MessageNone, flamegpu::MessageNone) { + FLAMEGPU->setVariable("b", FLAMEGPU->environment.getMacroProperty("int")); + return flamegpu::ALIVE; +} +)###"; + +TEST(DeviceMacroPropertyTest, RTC_add) { + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("a"); + agent.newVariable("b"); + AgentFunctionDescription& initFn = agent.newRTCFunction("init", Init_add_func); + AgentFunctionDescription& writeFn = agent.newRTCFunction("write", Write_add_func); + AgentFunctionDescription& readFn = agent.newRTCFunction("read", Read_add_func); + model.newLayer().addAgentFunction(initFn); + model.newLayer().addAgentFunction(writeFn); + model.newLayer().addAgentFunction(readFn); + AgentVector population(agent, 1); + population[0].setVariable("a", 12u); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + ASSERT_NO_THROW(cudaSimulation.getPopulationData(population)); + const unsigned int t_out = population.at(0).getVariable("b"); + ASSERT_EQ(13u, t_out); +} + +} // namespace +} // namespace flamegpu diff --git a/tests/test_cases/runtime/test_host_macro_property.cu b/tests/test_cases/runtime/test_host_macro_property.cu new file mode 100644 index 000000000..5a8a54e19 --- /dev/null +++ b/tests/test_cases/runtime/test_host_macro_property.cu @@ -0,0 +1,551 @@ +/** + * Tests of class: HostMacroProperty + * ReadTest: Test that HostMacroProperty can read data (written via agent fn's) + * WriteTest: Test that HostMacroProperty can write data (read via agent fn's) + * ZeroTest: Test that HostMacroProperty can zero data (read via agent fn's) + */ + +#include + +#include "flamegpu/flamegpu.h" + +#include "gtest/gtest.h" + +namespace flamegpu { +namespace { +const unsigned int TEST_DIMS[4] = {2, 3, 4, 5}; +FLAMEGPU_STEP_FUNCTION(HostRead) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + unsigned int a = 0; + for (unsigned int i = 0; i < TEST_DIMS[0]; ++i) { + for (unsigned int j = 0; j < TEST_DIMS[1]; ++j) { + for (unsigned int k = 0; k < TEST_DIMS[2]; ++k) { + for (unsigned int w = 0; w < TEST_DIMS[3]; ++w) { + ASSERT_EQ(t[i][j][k][w], a++); + } + } + } + } +} +FLAMEGPU_AGENT_FUNCTION(AgentWrite, MessageNone, MessageNone) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + const unsigned int i = FLAMEGPU->getVariable("i"); + const unsigned int j = FLAMEGPU->getVariable("j"); + const unsigned int k = FLAMEGPU->getVariable("k"); + const unsigned int w = FLAMEGPU->getVariable("w"); + t[i][j][k][w].exchange(FLAMEGPU->getVariable("a")); + return flamegpu::ALIVE; +} +TEST(HostMacroPropertyTest, ReadTest) { + // Fill MacroProperty with DeviceAPI + // Test values match expected value with HostAPI and test in-place + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("i"); + agent.newVariable("j"); + agent.newVariable("k"); + agent.newVariable("w"); + agent.newVariable("a"); + agent.newFunction("agentwrite", AgentWrite); + model.newLayer().addAgentFunction(AgentWrite); + model.newLayer().addHostFunction(HostRead); + const unsigned int total_agents = TEST_DIMS[0] * TEST_DIMS[1] * TEST_DIMS[2] * TEST_DIMS[3]; + AgentVector population(agent, total_agents); + unsigned int a = 0; + for (unsigned int i = 0; i < TEST_DIMS[0]; ++i) { + for (unsigned int j = 0; j < TEST_DIMS[1]; ++j) { + for (unsigned int k = 0; k < TEST_DIMS[2]; ++k) { + for (unsigned int w = 0; w < TEST_DIMS[3]; ++w) { + auto p = population[a]; + p.setVariable("a", a++); + p.setVariable("i", i); + p.setVariable("j", j); + p.setVariable("k", k); + p.setVariable("w", w); + } + } + } + } + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); +} +FLAMEGPU_STEP_FUNCTION(HostWrite) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + unsigned int a = 0; + for (unsigned int i = 0; i < TEST_DIMS[0]; ++i) { + for (unsigned int j = 0; j < TEST_DIMS[1]; ++j) { + for (unsigned int k = 0; k < TEST_DIMS[2]; ++k) { + for (unsigned int w = 0; w < TEST_DIMS[3]; ++w) { + t[i][j][k][w] = 1+a++; + } + } + } + } +} +FLAMEGPU_AGENT_FUNCTION(AgentRead, MessageNone, MessageNone) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + const unsigned int i = FLAMEGPU->getVariable("i"); + const unsigned int j = FLAMEGPU->getVariable("j"); + const unsigned int k = FLAMEGPU->getVariable("k"); + const unsigned int w = FLAMEGPU->getVariable("w"); + if (2 + t[i][j][k][w] == FLAMEGPU->getVariable("a")+1) { + FLAMEGPU->setVariable("a", 12); + } else { + FLAMEGPU->setVariable("a", 0); + } + return flamegpu::ALIVE; +} +TEST(HostMacroPropertyTest, WriteTest) { + // Fill MacroProperty with HostAPI + // Test values match expected value with DeviceAPI + // Write results back to agent variable, and check at the end + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + model.Environment().newMacroProperty("plusequal"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("i"); + agent.newVariable("j"); + agent.newVariable("k"); + agent.newVariable("w"); + agent.newVariable("a"); + agent.newFunction("agentread", AgentRead); + model.newLayer().addHostFunction(HostWrite); + model.newLayer().addAgentFunction(AgentRead); + const unsigned int total_agents = TEST_DIMS[0] * TEST_DIMS[1] * TEST_DIMS[2] * TEST_DIMS[3]; + AgentVector population(agent, total_agents); + unsigned int a = 0; + for (unsigned int i = 0; i < TEST_DIMS[0]; ++i) { + for (unsigned int j = 0; j < TEST_DIMS[1]; ++j) { + for (unsigned int k = 0; k < TEST_DIMS[2]; ++k) { + for (unsigned int w = 0; w < TEST_DIMS[3]; ++w) { + auto p = population[a]; + p.setVariable("a", 2+a++); + p.setVariable("i", i); + p.setVariable("j", j); + p.setVariable("k", k); + p.setVariable("w", w); + } + } + } + } + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + cudaSimulation.getPopulationData(population); + // Check results + unsigned int correct = 0; + for (auto p : population) { + correct += p.getVariable("a") == 12 ? 1 : 0; + } + ASSERT_EQ(correct, total_agents); +} +FLAMEGPU_STEP_FUNCTION(HostZero) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + t.zero(); +} +FLAMEGPU_AGENT_FUNCTION(AgentReadZero, MessageNone, MessageNone) { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + const unsigned int i = FLAMEGPU->getVariable("i"); + const unsigned int j = FLAMEGPU->getVariable("j"); + const unsigned int k = FLAMEGPU->getVariable("k"); + const unsigned int w = FLAMEGPU->getVariable("w"); + if (t[i][j][k][w] == 0) { + FLAMEGPU->setVariable("a", 1); + } else { + FLAMEGPU->setVariable("a", 0); + } + return flamegpu::ALIVE; +} +TEST(HostMacroPropertyTest, ZeroTest) { + // Fill MacroProperty with HostAPI + // Test values match expected value with DeviceAPI + // Write results back to agent variable, and check at the end + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("i"); + agent.newVariable("j"); + agent.newVariable("k"); + agent.newVariable("w"); + agent.newVariable("a"); + agent.newFunction("agentwrite", AgentWrite); + agent.newFunction("agentread", AgentReadZero); + model.newLayer().addAgentFunction(AgentWrite); + model.newLayer().addHostFunction(HostZero); + model.newLayer().addAgentFunction(AgentReadZero); + const unsigned int total_agents = TEST_DIMS[0] * TEST_DIMS[1] * TEST_DIMS[2] * TEST_DIMS[3]; + AgentVector population(agent, total_agents); + unsigned int a = 0; + for (unsigned int i = 0; i < TEST_DIMS[0]; ++i) { + for (unsigned int j = 0; j < TEST_DIMS[1]; ++j) { + for (unsigned int k = 0; k < TEST_DIMS[2]; ++k) { + for (unsigned int w = 0; w < TEST_DIMS[3]; ++w) { + auto p = population[a++]; + p.setVariable("i", i); + p.setVariable("j", j); + p.setVariable("k", k); + p.setVariable("w", w); + } + } + } + } + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_NO_THROW(cudaSimulation.simulate()); + cudaSimulation.getPopulationData(population); + // Check results + unsigned int correct = 0; + for (auto p : population) { + correct += p.getVariable("a") == 1 ? 1 : 0; + } + ASSERT_EQ(correct, total_agents); +} +FLAMEGPU_STEP_FUNCTION(HostArithmeticInit) { + FLAMEGPU->environment.getMacroProperty("int") = 10; + FLAMEGPU->environment.getMacroProperty("uint") = 10u; + FLAMEGPU->environment.getMacroProperty("int8") = 10; + FLAMEGPU->environment.getMacroProperty("uint8") = 10u; + FLAMEGPU->environment.getMacroProperty("int64") = 10; + FLAMEGPU->environment.getMacroProperty("uint64") = 10u; + FLAMEGPU->environment.getMacroProperty("id") = 10u; + FLAMEGPU->environment.getMacroProperty("float") = 10.0f; + FLAMEGPU->environment.getMacroProperty("double") = 10.0; +} +FLAMEGPU_STEP_FUNCTION(HostArithmetic) { + { + auto t = FLAMEGPU->environment.getMacroProperty("int"); + EXPECT_EQ(t++, 10); + EXPECT_EQ(++t, 12); + EXPECT_EQ(t--, 12); + EXPECT_EQ(--t, 10); + EXPECT_EQ(t / 5, 10 / 5); + EXPECT_EQ(t / 3, 10 / 3); + EXPECT_EQ(t / 3.0, 10 / 3.0); + EXPECT_EQ(t + 5, 10 + 5); + EXPECT_EQ(t + 3.0, 10 + 3.0); + EXPECT_EQ(t - 3, 10 - 3); + EXPECT_EQ(t - 3.0, 10 - 3.0); + EXPECT_EQ(t * 5, 10 * 5); + EXPECT_EQ(t * 3.0, 10 * 3.0); + EXPECT_EQ(t % 5, 10 % 5); + EXPECT_EQ(t % 3, 10 % 3); + EXPECT_EQ(t, 10); + t += 10; + EXPECT_EQ(t, 20); + t -= 10; + EXPECT_EQ(t, 10); + t *= 2; + EXPECT_EQ(t, 20); + t /= 2; + EXPECT_EQ(t, 10); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("uint"); + EXPECT_EQ(t++, 10u); + EXPECT_EQ(++t, 12u); + EXPECT_EQ(t--, 12u); + EXPECT_EQ(--t, 10u); + EXPECT_EQ(t / 5, 10u / 5); + EXPECT_EQ(t / 3, 10u / 3); + EXPECT_EQ(t / 3.0, 10u / 3.0); + EXPECT_EQ(t + 5, 10u + 5); + EXPECT_EQ(t + 3.0, 10u + 3.0); + EXPECT_EQ(t - 3, 10u - 3); + EXPECT_EQ(t - 3.0, 10u - 3.0); + EXPECT_EQ(t * 5, 10u * 5); + EXPECT_EQ(t * 3.0, 10u * 3.0); + EXPECT_EQ(t % 5, 10u % 5); + EXPECT_EQ(t % 3, 10u % 3); + EXPECT_EQ(t, 10u); + t += 10; + EXPECT_EQ(t, 20u); + t -= 10; + EXPECT_EQ(t, 10u); + t *= 2; + EXPECT_EQ(t, 20u); + t /= 2; + EXPECT_EQ(t, 10u); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("int8"); + EXPECT_EQ(t++, 10); + EXPECT_EQ(++t, 12); + EXPECT_EQ(t--, 12); + EXPECT_EQ(--t, 10); + EXPECT_EQ(t / 5, 10 / 5); + EXPECT_EQ(t / 3, 10 / 3); + EXPECT_EQ(t / 3.0, 10 / 3.0); + EXPECT_EQ(t + 5, 10 + 5); + EXPECT_EQ(t + 3.0, 10 + 3.0); + EXPECT_EQ(t - 3, 10 - 3); + EXPECT_EQ(t - 3.0, 10 - 3.0); + EXPECT_EQ(t * 5, 10 * 5); + EXPECT_EQ(t * 3.0, 10 * 3.0); + EXPECT_EQ(t % 5, 10 % 5); + EXPECT_EQ(t % 3, 10 % 3); + EXPECT_EQ(t, 10); + t += static_cast(10); + EXPECT_EQ(t, 20); + t -= static_cast(10); + EXPECT_EQ(t, 10); + t *= static_cast(2); + EXPECT_EQ(t, 20); + t /= static_cast(2); + EXPECT_EQ(t, 10); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("uint8"); + EXPECT_EQ(t++, 10u); + EXPECT_EQ(++t, 12u); + EXPECT_EQ(t--, 12u); + EXPECT_EQ(--t, 10u); + EXPECT_EQ(t / 5, static_cast(10u) / 5); + EXPECT_EQ(t / 3, static_cast(10u) / 3); + EXPECT_EQ(t / 3.0, static_cast(10u) / 3.0); + EXPECT_EQ(t + 5, static_cast(10u) + 5); + EXPECT_EQ(t + 3.0, static_cast(10u) + 3.0); + EXPECT_EQ(t - 3, static_cast(10u) - 3); + EXPECT_EQ(t - 3.0, static_cast(10u) - 3.0); + EXPECT_EQ(t * 5, static_cast(10u) * 5); + EXPECT_EQ(t * 3.0, static_cast(10u) * 3.0); + EXPECT_EQ(t % 5, static_cast(10u) % 5); + EXPECT_EQ(t % 3, static_cast(10u) % 3); + EXPECT_EQ(t, 10u); + t += static_cast(10u); + EXPECT_EQ(t, 20u); + t -= static_cast(10u); + EXPECT_EQ(t, 10u); + t *= static_cast(2u); + EXPECT_EQ(t, 20u); + t /= static_cast(2u); + EXPECT_EQ(t, 10u); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("int64"); + EXPECT_EQ(t++, 10); + EXPECT_EQ(++t, 12); + EXPECT_EQ(t--, 12); + EXPECT_EQ(--t, 10); + EXPECT_EQ(t / 5, 10 / 5); + EXPECT_EQ(t / 3, 10 / 3); + EXPECT_EQ(t / 3.0, 10 / 3.0); + EXPECT_EQ(t + 5, 10 + 5); + EXPECT_EQ(t + 3.0, 10 + 3.0); + EXPECT_EQ(t - 3, 10 - 3); + EXPECT_EQ(t - 3.0, 10 - 3.0); + EXPECT_EQ(t * 5, 10 * 5); + EXPECT_EQ(t * 3.0, 10 * 3.0); + EXPECT_EQ(t % 5, 10 % 5); + EXPECT_EQ(t % 3, 10 % 3); + EXPECT_EQ(t, 10); + t += 10; + EXPECT_EQ(t, 20); + t -= 10; + EXPECT_EQ(t, 10); + t *= 2; + EXPECT_EQ(t, 20); + t /= 2; + EXPECT_EQ(t, 10); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("uint64"); + EXPECT_EQ(t++, 10u); + EXPECT_EQ(++t, 12u); + EXPECT_EQ(t--, 12u); + EXPECT_EQ(--t, 10u); + EXPECT_EQ(t / 5, 10u / 5); + EXPECT_EQ(t / 3, 10u / 3); + EXPECT_EQ(t / 3.0, 10u / 3.0); + EXPECT_EQ(t + 5, 10u + 5); + EXPECT_EQ(t + 3.0, 10u + 3.0); + EXPECT_EQ(t - 3, 10u - 3); + EXPECT_EQ(t - 3.0, 10u - 3.0); + EXPECT_EQ(t * 5, 10u * 5); + EXPECT_EQ(t * 3.0, 10u * 3.0); + EXPECT_EQ(t % 5, 10u % 5); + EXPECT_EQ(t % 3, 10u % 3); + EXPECT_EQ(t, 10u); + t += 10; + EXPECT_EQ(t, 20u); + t -= 10; + EXPECT_EQ(t, 10u); + t *= 2; + EXPECT_EQ(t, 20u); + t /= 2; + EXPECT_EQ(t, 10u); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("id"); + EXPECT_EQ(t++, 10u); + EXPECT_EQ(++t, 12u); + EXPECT_EQ(t--, 12u); + EXPECT_EQ(--t, 10u); + EXPECT_EQ(t / 5, 10u / 5); + EXPECT_EQ(t / 3, 10u / 3); + EXPECT_EQ(t / 3.0, 10u / 3.0); + EXPECT_EQ(t + 5, 10u + 5); + EXPECT_EQ(t + 3.0, 10u + 3.0); + EXPECT_EQ(t - 3, 10u - 3); + EXPECT_EQ(t - 3.0, 10u - 3.0); + EXPECT_EQ(t * 5, 10u * 5); + EXPECT_EQ(t * 3.0, 10u * 3.0); + EXPECT_EQ(t % 5, 10u % 5); + EXPECT_EQ(t % 3, 10u % 3); + EXPECT_EQ(t, 10u); + t += 10; + EXPECT_EQ(t, 20u); + t -= 10; + EXPECT_EQ(t, 10u); + t *= 2; + EXPECT_EQ(t, 20u); + t /= 2; + EXPECT_EQ(t, 10u); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("float"); + EXPECT_EQ(t++, 10); + EXPECT_EQ(++t, 12); + EXPECT_EQ(t--, 12); + EXPECT_EQ(--t, 10); + EXPECT_EQ(t / 5, 10.0f / 5); + EXPECT_EQ(t / 3, 10.0f / 3); + EXPECT_EQ(t + 5, 10.0f + 5); + EXPECT_EQ(t - 3, 10.0f - 3); + EXPECT_EQ(t * 5, 10.0f * 5); + // EXPECT_EQ(t % 5, 10.0f % 5); // remainder op does not support floating point types + // EXPECT_EQ(t % 3, 10.0f % 3); // remainder op does not support floating point types + EXPECT_EQ(t, 10); + t += 10; + EXPECT_EQ(t, 20); + t -= 10; + EXPECT_EQ(t, 10); + t *= 2; + EXPECT_EQ(t, 20); + t /= 2; + EXPECT_EQ(t, 10); + } + { + auto t = FLAMEGPU->environment.getMacroProperty("double"); + EXPECT_EQ(t++, 10); + EXPECT_EQ(++t, 12); + EXPECT_EQ(t--, 12); + EXPECT_EQ(--t, 10); + EXPECT_EQ(t / 5, 10.0 / 5); + EXPECT_EQ(t / 3, 10.0 / 3); + EXPECT_EQ(t + 5, 10.0 + 5); + EXPECT_EQ(t - 3, 10.0 - 3); + EXPECT_EQ(t * 5, 10.0 * 5); + // EXPECT_EQ(t % 5, 10.0 % 5); // remainder op does not support floating point types + // EXPECT_EQ(t % 3, 10.0 % 3); // remainder op does not support floating point types + EXPECT_EQ(t, 10); + t += 10; + EXPECT_EQ(t, 20); + t -= 10; + EXPECT_EQ(t, 10); + t *= 2; + EXPECT_EQ(t, 20); + t /= 2; + EXPECT_EQ(t, 10); + } +} +TEST(HostMacroPropertyTest, ArithmeticTest) { + // Create single macro property for each type + // Fill MacroProperties with HostAPI to a known value + // Use all airthmetic ops, and test values match expected value with DeviceAPI + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + model.Environment().newMacroProperty("uint"); + model.Environment().newMacroProperty("int8"); + model.Environment().newMacroProperty("uint8"); + model.Environment().newMacroProperty("int64"); + model.Environment().newMacroProperty("uint64"); + model.Environment().newMacroProperty("id"); + model.Environment().newMacroProperty("float"); + model.Environment().newMacroProperty("double"); + // Setup agent fn + model.newAgent("agent"); + model.newLayer().addHostFunction(HostArithmeticInit); + model.newLayer().addHostFunction(HostArithmetic); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + ASSERT_NO_THROW(cudaSimulation.simulate()); +} + +/* These tests, test functionality which is not exposed unless LayerDescription allows agent fn and host fn in the same layer +#if !defined(SEATBELTS) || SEATBELTS +TEST(HostMacroPropertyTest, ReadSameLayerAsAgentWrite) { +#else +TEST(HostMacroPropertyTest, DISABLED_ReadSameLayerAsAgentWrite) { +#endif + // Ensure an exception is thrown if an agent fn writes to the same macro property as a host function reads (or writes) from in the same layer + // This is currently safe, as we do not execute host and agent functions concurrently, but may change in future + + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("i"); + agent.newVariable("j"); + agent.newVariable("k"); + agent.newVariable("w"); + agent.newVariable("a"); + agent.newFunction("agentwrite", AgentWrite); + auto &l = model.newLayer(); + l.addAgentFunction(AgentWrite); + l.addHostFunction(HostRead); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_THROW(cudaSimulation.simulate(), flamegpu::exception::InvalidOperation); +} +#if !defined(SEATBELTS) || SEATBELTS +TEST(HostMacroPropertyTest, WriteSameLayerAsAgentRead) { +#else +TEST(HostMacroPropertyTest, DISABLED_WriteSameLayerAsAgentRead) { +#endif + // Ensure an exception is thrown if an agent fn reads from the same macro property as a host function writes to in the same layer + // This is currently safe, as we do not execute host and agent functions concurrently, but may change in future + ModelDescription model("device_env_test"); + // Setup environment + model.Environment().newMacroProperty("int"); + model.Environment().newMacroProperty("plusequal"); + // Setup agent fn + AgentDescription& agent = model.newAgent("agent"); + agent.newVariable("i"); + agent.newVariable("j"); + agent.newVariable("k"); + agent.newVariable("w"); + agent.newVariable("a"); + agent.newFunction("agentread", AgentRead); + auto &l = model.newLayer(); + l.addHostFunction(HostWrite); + l.addAgentFunction(AgentRead); + AgentVector population(agent, 1); + // Do Sim + CUDASimulation cudaSimulation(model); + cudaSimulation.SimulationConfig().steps = 1; + cudaSimulation.setPopulationData(population); + ASSERT_THROW(cudaSimulation.simulate(), flamegpu::exception::InvalidOperation); +} +*/ +} // namespace +} // namespace flamegpu