.. _program_listing_file_include_flamegpu_runtime_messaging_MessageBruteForce_MessageBruteForceDevice.cuh: Program Listing for File MessageBruteForceDevice.cuh ==================================================== |exhale_lsh| :ref:`Return to documentation for file ` (``include/flamegpu/runtime/messaging/MessageBruteForce/MessageBruteForceDevice.cuh``) .. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS .. code-block:: cpp #ifndef INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_MESSAGEBRUTEFORCE_MESSAGEBRUTEFORCEDEVICE_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_MESSAGEBRUTEFORCE_MESSAGEBRUTEFORCEDEVICE_CUH_ #include "flamegpu/defines.h" #include "flamegpu/runtime/messaging/MessageNone.h" #include "flamegpu/runtime/messaging/MessageBruteForce.h" #ifndef __CUDACC_RTC__ #include "flamegpu/runtime/detail/curve/DeviceCurve.cuh" #endif // __CUDACC_RTC__ struct ModelData; namespace flamegpu { class MessageBruteForce::In { public: class Message; // Forward declare inner classes class iterator; // Forward declare inner classes __device__ In(const void *metadata) : len(reinterpret_cast(metadata)->length) { } __device__ size_type size(void) const { return len; } __device__ iterator begin(void) const { // const return iterator(*this, 0); } __device__ iterator end(void) const { // const // If there can be many begin, each with diff end, we need a middle layer to host the iterator/s return iterator(*this, len); } class Message { const MessageBruteForce::In &_parent; size_type index; public: __device__ Message(const MessageBruteForce::In &parent) : _parent(parent), index(0) {} __device__ Message(const MessageBruteForce::In &parent, size_type index) : _parent(parent), index(index) {} __host__ __device__ bool operator==(const Message& rhs) const { return this->getIndex() == rhs.getIndex(); } __host__ __device__ bool operator!=(const Message& rhs) const { return this->getIndex() != rhs.getIndex(); } __host__ __device__ Message& operator++() { ++index; return *this; } __host__ __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 iterator { Message _message; public: __device__ iterator(const In &parent, size_type index) : _message(parent, index) {} __device__ iterator& operator++() { ++_message; return *this; } __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; } }; private: size_type len; }; class MessageBruteForce::Out { public: __device__ Out(const void *, unsigned int *scan_flag_messageOutput) : scan_flag(scan_flag_messageOutput) { } 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; }; template __device__ T MessageBruteForce::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.len) { DTHROW("Brute force message index exceeds messagelist length, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. #ifdef FLAMEGPU_USE_GLM T value = detail::curve::DeviceCurve::getMessageVariable(variable_name, index); #else T value = detail::curve::DeviceCurve::getMessageVariable_ldg(variable_name, index); #endif return value; } template __device__ T MessageBruteForce::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) const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Ensure that the message is within bounds. if (index >= this->_parent.len) { DTHROW("Brute force message index exceeds messagelist length, unable to get variable '%s'.\n", variable_name); return static_cast(0); } #endif // get the value from curve using the message index. #ifdef FLAMEGPU_USE_GLM T value = detail::curve::DeviceCurve::getMessageArrayVariable(variable_name, index, array_index); #else T value = detail::curve::DeviceCurve::getMessageArrayVariable_ldg(variable_name, index, array_index); #endif return value; } template __device__ void MessageBruteForce::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 MessageBruteForce::Out::setVariable().\n", variable_name); #endif return; // Fail silently } unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // + d_message_count; // Todo: checking if the output message type is single or optional? (d_message_type) // set the variable using curve detail::curve::DeviceCurve::setMessageVariable(variable_name, value, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } template __device__ void MessageBruteForce::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 MessageBruteForce::Out::setVariable().\n", variable_name); #endif return; // Fail silently } unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // Todo: checking if the output message type is single or optional? (d_message_type) // set the variable using curve detail::curve::DeviceCurve::setMessageArrayVariable(variable_name, value, index, array_index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_MESSAGEBRUTEFORCE_MESSAGEBRUTEFORCEDEVICE_CUH_