Program Listing for File DeviceCurve.cuh

↰ Return to documentation for file (include/flamegpu/runtime/detail/curve/DeviceCurve.cuh)

#ifndef INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_DEVICECURVE_CUH_
#define INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_DEVICECURVE_CUH_

#include "flamegpu/runtime/detail/SharedBlock.h"
#include "flamegpu/runtime/detail/curve/Curve.cuh"
#include "flamegpu/exception/FLAMEGPUDeviceException_device.cuh"
#include "flamegpu/detail/type_decode.h"
#include "flamegpu/util/dstring.h"

#ifdef FLAMEGPU_USE_GLM
#ifdef __CUDACC__
#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
#pragma nv_diag_suppress = esa_on_defaulted_function_ignored
#else
#pragma diag_suppress = esa_on_defaulted_function_ignored
#endif  // __NVCC_DIAG_PRAGMA_SUPPORT__
#endif  // __CUDACC__
#include <glm/glm.hpp>
#endif  // FLAMEGPU_USE_GLM

namespace flamegpu {
namespace detail {
namespace curve {

class DeviceCurve {
 public:
    typedef Curve::Variable          Variable;                    // !< Typedef for cuRVE variable handle
    typedef Curve::VariableHash      VariableHash;                // !< Typedef for cuRVE variable name string hash
    typedef Curve::NamespaceHash     NamespaceHash;               // !< Typedef for cuRVE variable namespace string hash
    static const int MAX_VARIABLES = Curve::MAX_VARIABLES;        // !< Default maximum number of cuRVE variables (must be a power of 2)static const VariableHash EMPTY_FLAG = Curve::EMPTY_FLAG;
    static const VariableHash EMPTY_FLAG = Curve::EMPTY_FLAG;
    static const int UNKNOWN_VARIABLE = -1;                       // !< value returned as a Variable if an API function encounters an error


    __device__ __forceinline__ static Variable getVariableIndex(VariableHash variable_hash);

 private:
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static char* getVariablePtr(const char(&variableName)[M], VariableHash namespace_hash, unsigned int offset);

    template <typename T, unsigned int N = 1, unsigned int M>
    __device__ __forceinline__ static T getVariable(const char(&variableName)[M], VariableHash namespace_hash, unsigned int agent_index = 0, unsigned int array_index = 0);
    template <typename T, unsigned int N = 1, unsigned int M>
    __device__ __forceinline__ static T getVariable_ldg(const char(&variableName)[M], VariableHash namespace_hash, unsigned int agent_index = 0, unsigned int array_index = 0);
    template <typename T, unsigned int N = 1, unsigned int M>
    __device__ __forceinline__ static void setVariable(const char(&variableName)[M], VariableHash namespace_hash, T value, unsigned int agent_index = 0, unsigned int array_index = 0);

 public:
     __device__ __forceinline__ static void init(const CurveTable* __restrict__ d_curve_table) {
         using detail::sm;
         for (int idx = threadIdx.x; idx < Curve::MAX_VARIABLES; idx += blockDim.x) {
             sm()->curve_variables[idx] = d_curve_table->variables[idx];
             sm()->curve_hashes[idx] = d_curve_table->hashes[idx];
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
             sm()->curve_type_size[idx] = d_curve_table->type_size[idx];
             sm()->curve_elements[idx] = d_curve_table->elements[idx];
#endif
             sm()->curve_count[idx] = d_curve_table->count[idx];
         }
    }

    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getAgentVariable(const char(&variableName)[M], unsigned int index);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getMessageVariable(const char(&variableName)[M], unsigned int index);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getAgentVariable_ldg(const char(&variableName)[M], unsigned int index);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getMessageVariable_ldg(const char(&variableName)[M], unsigned int index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static T getAgentArrayVariable(const char(&variableName)[M], unsigned int variable_index, unsigned int array_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static T getMessageArrayVariable(const char(&variableName)[M], unsigned int variable_index, unsigned int array_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static T getAgentArrayVariable_ldg(const char(&variableName)[M], unsigned int variable_index, unsigned int array_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static T getMessageArrayVariable_ldg(const char(&variableName)[M], unsigned int variable_index, unsigned int array_index);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static void setAgentVariable(const char(&variableName)[M], T variable, unsigned int index);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static void setMessageVariable(const char(&variableName)[M], T variable, unsigned int index);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static void setNewAgentVariable(const char(&variableName)[M], T variable, unsigned int index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static void setAgentArrayVariable(const char(&variableName)[M], T variable, unsigned int variable_index, unsigned int array_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static void setMessageArrayVariable(const char(&variableName)[M], T variable, unsigned int variable_index, unsigned int array_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static void setNewAgentArrayVariable(const char(&variableName)[M], T variable, unsigned int variable_index, unsigned int array_index);

    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getEnvironmentProperty(const char(&propertyName)[M]);
    template <typename T, unsigned int N = 0, unsigned int M>
    __device__ __forceinline__ static T getEnvironmentArrayProperty(const char(&propertyName)[M], unsigned int array_index);


    __device__ __forceinline__ static unsigned int *getEnvironmentDirectedGraphPBM(VariableHash graphHash);
    __device__ __forceinline__ static unsigned int* getEnvironmentDirectedGraphIPBM(VariableHash graphHash);
    __device__ __forceinline__ static unsigned int* getEnvironmentDirectedGraphIPBMEdges(VariableHash graphHash);
    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getEnvironmentDirectedGraphVertexProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int vertex_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static T getEnvironmentDirectedGraphVertexArrayProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int vertex_index, unsigned int array_index);

    template <typename T, unsigned int M>
    __device__ __forceinline__ static T getEnvironmentDirectedGraphEdgeProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int edge_index);
    template <typename T, unsigned int N, unsigned int M>
    __device__ __forceinline__ static T getEnvironmentDirectedGraphEdgeArrayProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int edge_index, unsigned int array_index);

    template<typename T, unsigned int I = 1, unsigned int J = 1, unsigned int K = 1, unsigned int W = 1, unsigned int M>
    __device__ __forceinline__ static char *getEnvironmentMacroProperty(const char(&name)[M]);

    template <unsigned int M>
    __device__ __forceinline__ static unsigned int getVariableCount(const char(&variableName)[M], VariableHash namespace_hash);

    __device__ __forceinline__ static bool isAgent(const char* agent_name);
    __device__ __forceinline__ static bool isState(const char* agent_state);
};

__device__ __forceinline__ DeviceCurve::Variable DeviceCurve::getVariableIndex(const VariableHash variable_hash) {
    using detail::sm;
    // loop unrolling of hash collision detection
    // (This may inflate register usage based on the max number of iterations in some cases)
    for (unsigned int x = 0; x< MAX_VARIABLES; x++) {
        const Variable i = (variable_hash + x) & (MAX_VARIABLES - 1);
        if (sm()->curve_hashes[i] == variable_hash)
            return i;
    }
    return UNKNOWN_VARIABLE;
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ char* DeviceCurve::getVariablePtr(const char(&variableName)[M], const VariableHash namespace_hash, const unsigned int offset) {
    using detail::sm;
    const Variable cv = getVariableIndex(Curve::variableHash(variableName) + namespace_hash);
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (cv == UNKNOWN_VARIABLE) {
        DTHROW("Curve variable with name '%s' was not found.\n", variableName);
        return nullptr;
    } else if (sm()->curve_type_size[cv] != sizeof(typename detail::type_decode<T>::type_t)) {
        DTHROW("Curve variable with name '%s', type size mismatch %u != %llu.\n", variableName, sm()->curve_type_size[cv], sizeof(typename detail::type_decode<T>::type_t));
        return nullptr;
    } else if (!(sm()->curve_elements[cv] == detail::type_decode<T>::len_t * N || (namespace_hash == Curve::variableHash("_environment") && N == 0))) {  // Special case, environment can avoid specifying N
        DTHROW("Curve variable with name '%s', variable array length mismatch %u != %u.\n", variableName, sm()->curve_elements[cv], detail::type_decode<T>::len_t);
        return nullptr;
    } else if (offset >= sm()->curve_type_size[cv] * sm()->curve_elements[cv] * sm()->curve_count[cv]) {  // Note : offset is basically index * sizeof(T)
        DTHROW("Curve variable with name '%s', offset exceeds buffer length  %u >= %u.\n", variableName, offset, sm()->curve_type_size[cv] * sm()->curve_elements[cv] * sm()->curve_count[cv]);
        return nullptr;
    }
#endif
    // return a generic pointer to variable address for given offset
    return sm()->curve_variables[cv] + offset;
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getVariable(const char(&variableName)[M], const VariableHash namespace_hash, const unsigned int agent_index, const unsigned int array_index) {
    using detail::sm;
    const unsigned int buffer_offset = agent_index * static_cast<unsigned int>(sizeof(T)) * N + array_index * sizeof(typename detail::type_decode<T>::type_t);
    T *value_ptr = reinterpret_cast<T*>(getVariablePtr<T, N>(variableName, namespace_hash, buffer_offset));

#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (!value_ptr)
        return {};
#endif
    return *value_ptr;
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getVariable_ldg(const char(&variableName)[M], const VariableHash namespace_hash, const unsigned int agent_index, const unsigned int array_index) {
    const unsigned int buffer_offset = agent_index * static_cast<unsigned int>(sizeof(T)) * N + array_index * sizeof(typename detail::type_decode<T>::type_t);
    T *value_ptr = reinterpret_cast<T*>(getVariablePtr<T, N>(variableName, namespace_hash, buffer_offset));

#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (!value_ptr)
        return {};
#endif
    return __ldg(value_ptr);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setVariable(const char(&variableName)[M], const VariableHash namespace_hash, const T variable, const unsigned int agent_index, const unsigned int array_index) {
    const unsigned int buffer_offset = agent_index * static_cast<unsigned int>(sizeof(T)) * N + array_index * sizeof(typename detail::type_decode<T>::type_t);
    T* value_ptr = reinterpret_cast<T*>(getVariablePtr<T, N>(variableName, namespace_hash, buffer_offset));

#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (!value_ptr)
        return;
#endif
    *value_ptr = variable;
}
template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getAgentVariable(const char (&variableName)[M], unsigned int index) {
    using detail::sm;
    return getVariable<T>(variableName, 0, index);
}
template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getMessageVariable(const char (&variableName)[M], unsigned int index) {
    return getVariable<T>(variableName, Curve::variableHash("_message_in"), index);
}
template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getAgentVariable_ldg(const char (&variableName)[M], unsigned int index) {
    return getVariable_ldg<T>(variableName, 0, index);
}
template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getMessageVariable_ldg(const char (&variableName)[M], unsigned int index) {
    return getVariable_ldg<T>(variableName, Curve::variableHash("_message_in"), index);
}

template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getAgentArrayVariable(const char(&variableName)[M], unsigned int agent_index, unsigned int array_index) {
    return getVariable<T, N>(variableName, 0, agent_index, array_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getMessageArrayVariable(const char(&variableName)[M], unsigned int message_index, unsigned int array_index) {
    return getVariable<T, N>(variableName, Curve::variableHash("_message_in"), message_index, array_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getAgentArrayVariable_ldg(const char(&variableName)[M], unsigned int agent_index, unsigned int array_index) {
    return getVariable_ldg<T, N>(variableName, 0, agent_index, array_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getMessageArrayVariable_ldg(const char(&variableName)[M], unsigned int message_index, unsigned int array_index) {
    return getVariable_ldg<T, N>(variableName, Curve::variableHash("_message_in"), message_index, array_index);
}

template <typename T, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setAgentVariable(const char(&variableName)[M], T variable, unsigned int index) {
    setVariable<T>(variableName, 0, variable, index);
}
template <typename T, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setMessageVariable(const char(&variableName)[M], T variable, unsigned int index) {
    setVariable<T>(variableName, Curve::variableHash("_message_out"), variable, index);
}
template <typename T, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setNewAgentVariable(const char(&variableName)[M], T variable, unsigned int index) {
    setVariable<T>(variableName, Curve::variableHash("_agent_birth"), variable, index);
}

template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setAgentArrayVariable(const char(&variableName)[M], T variable, unsigned int agent_index, unsigned int array_index) {
    setVariable<T, N>(variableName, 0, variable, agent_index, array_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setMessageArrayVariable(const char(&variableName)[M], T variable, unsigned int message_index, unsigned int array_index) {
    setVariable<T, N>(variableName, Curve::variableHash("_message_out"), variable, message_index, array_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void DeviceCurve::setNewAgentArrayVariable(const char(&variableName)[M], T variable, unsigned int agent_index, unsigned int array_index) {
    setVariable<T, N>(variableName, Curve::variableHash("_agent_birth"), variable, agent_index, array_index);
}

template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getEnvironmentProperty(const char(&propertyName)[M]) {
    using detail::sm;
    return  *reinterpret_cast<const T*>(sm()->env_buffer + reinterpret_cast<ptrdiff_t>(getVariablePtr<T, 1>(propertyName, Curve::variableHash("_environment"), 0)));
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getEnvironmentArrayProperty(const char(&propertyName)[M], unsigned int array_index) {
    using detail::sm;
    return *reinterpret_cast<const T*>(sm()->env_buffer + reinterpret_cast<ptrdiff_t>(getVariablePtr<T, N>(propertyName, Curve::variableHash("_environment"), array_index * sizeof(T))));
}
__device__ __forceinline__ unsigned int *DeviceCurve::getEnvironmentDirectedGraphPBM(VariableHash graphHash) {
    using detail::sm;
    return reinterpret_cast<unsigned int *>(getVariablePtr<unsigned int, 1>("_pbm", graphHash ^ Curve::variableHash("_environment_directed_graph_vertex"), 0));
}
__device__ __forceinline__ unsigned int* DeviceCurve::getEnvironmentDirectedGraphIPBM(VariableHash graphHash) {
    using detail::sm;
    return reinterpret_cast<unsigned int*>(getVariablePtr<unsigned int, 1>("_ipbm", graphHash ^ Curve::variableHash("_environment_directed_graph_vertex"), 0));
}
__device__ __forceinline__ unsigned int* DeviceCurve::getEnvironmentDirectedGraphIPBMEdges(VariableHash graphHash) {
    using detail::sm;
    return reinterpret_cast<unsigned int*>(getVariablePtr<unsigned int, 1>("_ipbm_edges", graphHash ^ Curve::variableHash("_environment_directed_graph_vertex"), 0));
}
template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getEnvironmentDirectedGraphVertexProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int vertex_index) {
    using detail::sm;
    return getVariable<T>(propertyName, graphHash ^ Curve::variableHash("_environment_directed_graph_vertex"), vertex_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getEnvironmentDirectedGraphVertexArrayProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int vertex_index, unsigned int array_index) {
    using detail::sm;
    return getVariable<T, N>(propertyName, graphHash ^ Curve::variableHash("_environment_directed_graph_vertex"), vertex_index, array_index);
}
template <typename T, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getEnvironmentDirectedGraphEdgeProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int edge_index) {
    using detail::sm;
    return getVariable<T>(propertyName, graphHash ^ Curve::variableHash("_environment_directed_graph_edge"), edge_index);
}
template <typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T DeviceCurve::getEnvironmentDirectedGraphEdgeArrayProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int edge_index, unsigned int array_index) {
    using detail::sm;
    return getVariable<T, N>(propertyName, graphHash ^ Curve::variableHash("_environment_directed_graph_edge"), edge_index, array_index);
}

template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W, unsigned int M>
__device__ __forceinline__ char* DeviceCurve::getEnvironmentMacroProperty(const char(&name)[M]) {
    return getVariablePtr<T, I*J*K*W>(name, Curve::variableHash("_macro_environment"), 0);
}

__device__ __forceinline__ bool DeviceCurve::isAgent(const char* agent_name) {
    return util::dstrcmp(agent_name, sm()->agent_name) == 0;
}
__device__ __forceinline__ bool DeviceCurve::isState(const char* agent_state) {
    return util::dstrcmp(agent_state, sm()->state_name) == 0;
}

template <unsigned int M>
__device__ __forceinline__ unsigned int DeviceCurve::getVariableCount(const char(&variableName)[M], const VariableHash namespace_hash) {
    using detail::sm;
    const Variable cv = getVariableIndex(Curve::variableHash(variableName) + namespace_hash);
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (cv == UNKNOWN_VARIABLE) {
        DTHROW("Curve variable with name '%s' was not found.\n", variableName);
        return 0;
    }
#endif
    return sm()->curve_count[cv];
}
}  // namespace curve
}  // namespace detail
}  // namespace flamegpu

#endif  // INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_DEVICECURVE_CUH_