.. _program_listing_file_include_flamegpu_runtime_environment_HostMacroProperty.cuh: Program Listing for File HostMacroProperty.cuh ============================================== |exhale_lsh| :ref:`Return to documentation for file ` (``include/flamegpu/runtime/environment/HostMacroProperty.cuh``) .. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS .. code-block:: cpp #ifndef INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_HOSTMACROPROPERTY_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_HOSTMACROPROPERTY_CUH_ #include #include #include #include #include #include namespace flamegpu { struct HostMacroProperty_MetaData { HostMacroProperty_MetaData(void* _d_base_ptr, const std::array& _dims, size_t _type_size, bool _device_read_flag, const std::string &name, cudaStream_t _stream) : 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) , stream(_stream) { } ~HostMacroProperty_MetaData() { upload(); if (h_base_ptr) std::free(h_base_ptr); } void download() { if (!h_base_ptr) { force_download(); } } void force_download() { if (!h_base_ptr) { h_base_ptr = static_cast(malloc(elements * type_size)); } gpuErrchk(cudaMemcpyAsync(h_base_ptr, d_base_ptr, elements * type_size, cudaMemcpyDeviceToHost, stream)); gpuErrchk(cudaStreamSynchronize(stream)); has_changed = false; } void upload() { if (h_base_ptr && has_changed) { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_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(cudaMemcpyAsync(d_base_ptr, h_base_ptr, elements * type_size, cudaMemcpyHostToDevice, stream)); gpuErrchk(cudaStreamSynchronize(stream)); 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; cudaStream_t stream; }; template class HostMacroProperty { ptrdiff_t offset; std::shared_ptr metadata; public: explicit HostMacroProperty(const std::shared_ptr& _metadata); HostMacroProperty(const std::shared_ptr &_metadata, const ptrdiff_t& _offset); HostMacroProperty operator[](unsigned int i) const; operator T(); operator T() const; HostMacroProperty& operator=(T val); void zero(); HostMacroProperty& operator++(); T operator++(int); HostMacroProperty& operator--(); T operator--(int); // Swig breaks if it see's auto return type #ifndef SWIG template auto operator+(T2 b) const; template auto operator-(T2 b) const; template auto operator*(T2 b) const; template auto operator/(T2 b) const; template auto operator%(T2 b) const; #endif // HostMacroProperty& operator=(T b); // Defined above template HostMacroProperty& operator+=(T2 b); template HostMacroProperty& operator-=(T2 b); template HostMacroProperty& operator*=(T2 b); template HostMacroProperty& operator/=(T2 b); template HostMacroProperty& operator%=(T2 b); private: T& _get(); T& _get() const; }; #ifdef SWIG template class HostMacroProperty_swig { const std::array dimensions; ptrdiff_t offset; std::shared_ptr metadata; public: explicit HostMacroProperty_swig(const std::shared_ptr &_metadata); HostMacroProperty_swig(const std::shared_ptr& _metadata, const ptrdiff_t& _offset, const std::array& _dimensions); HostMacroProperty_swig __getitem__(unsigned int i) const; void __setitem__(unsigned int i, T val); void set(T val); int __int__(); int64_t __long__(); double __float__(); bool __bool__(); bool __eq__(T other) const; bool __ne__(T other) const; bool __lt__(T other) const; bool __le__(T other) const; bool __gt__(T other) const; bool __ge__(T other) const; unsigned int __len__() const { return dimensions[0]; } T get() const; 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(FLAMEGPU_SEATBELTS) || FLAMEGPU_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_ENVIRONMENT_HOSTMACROPROPERTY_CUH_