.. _program_listing_file_include_flamegpu_runtime_messaging_MessageArray_MessageArrayDevice.cuh: Program Listing for File MessageArrayDevice.cuh =============================================== |exhale_lsh| :ref:`Return to documentation for file ` (``include/flamegpu/runtime/messaging/MessageArray/MessageArrayDevice.cuh``) .. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS .. code-block:: cpp #ifndef INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_MESSAGEARRAY_MESSAGEARRAYDEVICE_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_MESSAGEARRAY_MESSAGEARRAYDEVICE_CUH_ #include "flamegpu/runtime/messaging/MessageArray.h" #include "flamegpu/runtime/messaging/MessageBruteForce/MessageBruteForceDevice.cuh" namespace flamegpu { class MessageArray::In { friend class Message; public: class Message { const MessageArray::In &_parent; size_type index; public: __device__ Message(const MessageArray::In &parent, const size_type _index) : _parent(parent), index(_index) {} #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS __device__ Message(const MessageArray::In& parent) : _parent(parent), index(0) {} #endif __device__ bool operator==(const Message& rhs) const { return this->index == rhs.index; } __device__ bool operator!=(const Message& rhs) const { return this->index != rhs.index; } __device__ size_type getIndex() const { return this->index; } template __device__ T getVariable(const char(&variable_name)[N]) const; template __device__ T getVariable(const char(&variable_name)[M], unsigned int index) const; }; class WrapFilter { friend class Message; public: class Message { const WrapFilter&_parent; int relative_cell; size_type index_1d = 0; public: __device__ Message(const WrapFilter&parent, const int relative_x) : _parent(parent) { relative_cell = relative_x; } __device__ bool operator==(const Message& rhs) const { return this->relative_cell == rhs.relative_cell; // && this->_parent.loc == rhs._parent.loc; } __device__ bool operator!=(const Message& rhs) const { return !(*this == rhs); } __device__ inline Message& operator++(); __device__ size_type getX() const { return (this->_parent.loc + relative_cell + this->_parent.length) % this->_parent.length; } __device__ size_type getOffsetX() const { return relative_cell; } template __device__ T getVariable(const char(&variable_name)[N]) const; template __device__ T getVariable(const char(&variable_name)[M], unsigned int index) const; }; class iterator { Message _message; public: __device__ iterator(const WrapFilter&parent, const int relative_x) : _message(parent, relative_x) { // Increment to find first message ++_message; } __device__ iterator& operator++() { ++_message; return *this; } __device__ iterator operator++(int) { iterator temp = *this; ++*this; return temp; } __device__ bool operator==(const iterator& rhs) const { return _message == rhs._message; } __device__ bool operator!=(const iterator& rhs) const { return _message != rhs._message; } __device__ Message& operator*() { return _message; } __device__ Message* operator->() { return &_message; } }; inline __device__ WrapFilter(const size_type _length, const size_type x, const size_type _radius); #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS inline __device__ WrapFilter(); #endif inline __device__ iterator begin(void) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (!this->length) return iterator(*this, radius); #endif // Bin before initial bin, as the constructor calls increment operator return iterator(*this, -static_cast(radius) - 1); } inline __device__ iterator end(void) const { // Final bin, as the constructor calls increment operator return iterator(*this, radius); } private: size_type loc; const size_type radius; const size_type length; }; class Filter { friend class Message; public: class Message { const Filter& _parent; int relative_cell; size_type index_1d = 0; public: __device__ Message(const Filter& parent, const int relative_x) : _parent(parent) { relative_cell = relative_x; } __device__ bool operator==(const Message& rhs) const { return this->relative_cell == rhs.relative_cell; // && this->_parent.loc == rhs._parent.loc; } __device__ bool operator!=(const Message& rhs) const { return !(*this == rhs); } __device__ inline Message& operator++(); __device__ size_type getX() const { return this->_parent.loc + relative_cell; } __device__ int getOffsetX() const { return this->relative_cell; } template __device__ T getVariable(const char(&variable_name)[N]) const; template __device__ T getVariable(const char(&variable_name)[M], unsigned int index) const; }; class iterator { Message _message; public: __device__ iterator(const Filter& parent, const int relative_x) : _message(parent, relative_x) { // Increment to find first message ++_message; } __device__ iterator& operator++() { ++_message; return *this; } __device__ iterator operator++(int) { iterator temp = *this; ++* this; return temp; } __device__ bool operator==(const iterator& rhs) const { return _message == rhs._message; } __device__ bool operator!=(const iterator& rhs) const { return _message != rhs._message; } __device__ Message& operator*() { return _message; } __device__ Message* operator->() { return &_message; } }; inline __device__ Filter(size_type _length, size_type x, size_type _radius); #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS inline __device__ Filter(); #endif inline __device__ iterator begin(void) const { // Bin before initial bin, as the constructor calls increment operator return iterator(*this, min_cell - 1); } inline __device__ iterator end(void) const { // Final bin, as the constructor calls increment operator return iterator(*this, max_cell); } private: size_type loc; int min_cell; int max_cell; const size_type length; }; __device__ In(const void *metadata) : length(reinterpret_cast(metadata)->length) { } inline __device__ WrapFilter wrap(const size_type x, const size_type radius = 1) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (radius == 0) { DTHROW("Invalid radius %u for accessing array messagelist of length %u\n", radius, length); return WrapFilter(); } else if ((radius * 2) + 1 > length) { unsigned int min_r = length % 2 == 0 ? length - 2 : length - 1; min_r /= 2; DTHROW("%u is not a valid radius for accessing Array message lists, as the diameter of messages accessed exceeds the message list length (%u)." " Maximum supported radius for this message list is %u.\n", radius, length, min_r); return WrapFilter(); } else if (x >= length) { DTHROW("%u is not a valid position for iterating an Array message list of length %u, location must be within bounds.", x, length); return WrapFilter(); } #endif return WrapFilter(length, x, radius); } inline __device__ Filter operator() (const size_type x, const size_type radius = 1) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (radius == 0) { DTHROW("Invalid radius %u for accessing array messagelist of length %u\n", radius, length); return Filter(); } else if (x >= length) { DTHROW("%u is not a valid position for iterating an Array message list of length %u, location must be within bounds.", x, length); return Filter(); } #endif return Filter(length, x, radius); } __device__ size_type size(void) const { return length; } __device__ Message at(const size_type index) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (index >= length) { DTHROW("Index is out of bounds for Array messagelist (%u >= %u).\n", index, length); return Message(*this); } #endif return Message(*this, index); } private: const size_type length; }; class MessageArray::Out { public: __device__ Out(const void *_metadata, unsigned int *scan_flag_messageOutput) : scan_flag(scan_flag_messageOutput) #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS , metadata(reinterpret_cast(_metadata)) #else , metadata(nullptr) #endif { } __device__ inline void setIndex(const size_type id) const; template __device__ void setVariable(const char(&variable_name)[N], T value) const; template __device__ void setVariable(const char(&variable_name)[M], unsigned int index, T value) const; protected: unsigned int *scan_flag; const MetaData * const metadata; }; template __device__ T MessageArray::In::Message::getVariable(const char(&variable_name)[N]) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index >= this->_parent.length) { DTHROW("Invalid Array message, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. return detail::curve::DeviceCurve::getMessageVariable(variable_name, index); } template __device__ T MessageArray::In::Message::getVariable(const char(&variable_name)[M], const unsigned int array_index) const { // simple indexing assumes index is the thread number (this may change later) #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index >= this->_parent.length) { DTHROW("Invalid Array message, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. T value = detail::curve::DeviceCurve::getMessageArrayVariable(variable_name, index, array_index); return value; } template __device__ T MessageArray::In::WrapFilter::Message::getVariable(const char(&variable_name)[N]) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index_1d >= this->_parent.length) { DTHROW("Invalid Array message, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. return detail::curve::DeviceCurve::getMessageVariable(variable_name, index_1d); } template __device__ T MessageArray::In::WrapFilter::Message::getVariable(const char(&variable_name)[M], const unsigned int array_index) const { // simple indexing assumes index is the thread number (this may change later) #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index_1d >= this->_parent.length) { DTHROW("Invalid Array message, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. T value = detail::curve::DeviceCurve::getMessageArrayVariable(variable_name, index_1d, array_index); return value; } template __device__ T MessageArray::In::Filter::Message::getVariable(const char(&variable_name)[N]) const { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index_1d >= this->_parent.length) { DTHROW("Invalid Array message, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. return detail::curve::DeviceCurve::getMessageVariable(variable_name, index_1d); } template __device__ T MessageArray::In::Filter::Message::getVariable(const char(&variable_name)[M], const unsigned int array_index) const { // simple indexing assumes index is the thread number (this may change later) #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index_1d >= this->_parent.length) { DTHROW("Invalid Array message, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. T value = detail::curve::DeviceCurve::getMessageArrayVariable(variable_name, index_1d, array_index); return value; } template __device__ void MessageArray::Out::setVariable(const char(&variable_name)[N], T value) const { // message name or variable name if (variable_name[0] == '_') { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS DTHROW("Variable names starting with '_' are reserved for internal use, with '%s', in MessageArray::Out::setVariable().\n", variable_name); #endif return; // Fail silently } unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve detail::curve::DeviceCurve::setMessageVariable(variable_name, value, index); // setIndex() sets the optional message scan flag } template __device__ void MessageArray::Out::setVariable(const char(&variable_name)[M], const unsigned int array_index, T value) const { if (variable_name[0] == '_') { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS DTHROW("Variable names starting with '_' are reserved for internal use, with '%s', in MessageArray::Out::setVariable().\n", variable_name); #endif return; // Fail silently } unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve detail::curve::DeviceCurve::setMessageArrayVariable(variable_name, value, index, array_index); // setIndex() sets the optional message scan flag } __device__ void MessageArray::Out::setIndex(const size_type id) const { unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (id >= metadata->length) { DTHROW("MessageArray index [%u] is out of bounds [%u]\n", id, metadata->length); } #endif // set the variable using curve detail::curve::DeviceCurve::setMessageVariable("___INDEX", id, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } __device__ MessageArray::In::WrapFilter::WrapFilter(const size_type _length, const size_type x, const size_type _radius) : radius(_radius) , length(_length) { loc = x; } #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS __device__ inline MessageArray::In::WrapFilter::WrapFilter() : radius(0) , length(0) { loc = 0; } #endif __device__ MessageArray::In::WrapFilter::Message& MessageArray::In::WrapFilter::Message::operator++() { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (!_parent.length) return *this; #endif relative_cell++; // Skip origin cell if (relative_cell == 0) { relative_cell++; } // Wrap over boundaries index_1d = (this->_parent.loc + relative_cell + this->_parent.length) % this->_parent.length; return *this; } __device__ MessageArray::In::Filter::Filter(const size_type _length, const size_type x, const size_type _radius) : length(_length) { loc = x; min_cell = static_cast(x) - static_cast(_radius) < 0 ? -static_cast(x) : -static_cast(_radius); max_cell = x + _radius >= _length ? static_cast(_length) - 1 - static_cast(x) : static_cast(_radius); } #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS __device__ inline MessageArray::In::Filter::Filter() : length(0) { loc = 0; min_cell = 1; max_cell = 0; } #endif __device__ MessageArray::In::Filter::Message& MessageArray::In::Filter::Message::operator++() { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (!_parent.length) return *this; #endif relative_cell++; // Skip origin cell if (relative_cell == 0) { relative_cell++; } // Solve to 1 dimensional bin index index_1d = this->_parent.loc + relative_cell; return *this; } } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_MESSAGEARRAY_MESSAGEARRAYDEVICE_CUH_