.. _program_listing_file_include_flamegpu_runtime_environment_DeviceEnvironment.cuh: Program Listing for File DeviceEnvironment.cuh ============================================== |exhale_lsh| :ref:`Return to documentation for file ` (``include/flamegpu/runtime/environment/DeviceEnvironment.cuh``) .. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS .. code-block:: cpp #ifndef INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_DEVICEENVIRONMENT_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_DEVICEENVIRONMENT_CUH_ // #include #include #include #include "flamegpu/runtime/environment/DeviceMacroProperty.cuh" #include "flamegpu/runtime/environment/DeviceEnvironmentDirectedGraph.cuh" #include "flamegpu/detail/type_decode.h" #ifndef __CUDACC_RTC__ #include "flamegpu/runtime/detail/curve/DeviceCurve.cuh" #endif namespace flamegpu { class ReadOnlyDeviceEnvironment { friend class ReadOnlyDeviceAPI; public: template __device__ __forceinline__ T getProperty(const char(&name)[M]) const; template __device__ __forceinline__ T getProperty(const char(&name)[M], unsigned int index) const; template __device__ __forceinline__ ReadOnlyDeviceMacroProperty getMacroProperty(const char(&name)[M]) const; template __device__ __forceinline__ DeviceEnvironmentDirectedGraph getDirectedGraph(const char(&name)[M]) const; }; class DeviceEnvironment : public ReadOnlyDeviceEnvironment { public: template __device__ __forceinline__ DeviceMacroProperty getMacroProperty(const char(&name)[M]) const; }; // Mash compilation of these functions from RTC builds as this requires a dynamic implementation of the function in curve_rtc #ifndef __CUDACC_RTC__ template __device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[M]) const { return detail::curve::DeviceCurve::getEnvironmentProperty(name); } template __device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[M], const unsigned int index) const { return detail::curve::DeviceCurve::getEnvironmentArrayProperty(name, index); } template __device__ __forceinline__ ReadOnlyDeviceMacroProperty ReadOnlyDeviceEnvironment::getMacroProperty(const char(&name)[N]) const { char * d_ptr = detail::curve::DeviceCurve::getEnvironmentMacroProperty(name); #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (!d_ptr) { return ReadOnlyDeviceMacroProperty(nullptr, nullptr); } return ReadOnlyDeviceMacroProperty(reinterpret_cast(d_ptr), reinterpret_cast(d_ptr + (I * J * K * W * sizeof(T)))); // Read-write flag resides in 8 bits at the end of the buffer #else return ReadOnlyDeviceMacroProperty(reinterpret_cast(d_ptr)); #endif } template __device__ __forceinline__ DeviceMacroProperty DeviceEnvironment::getMacroProperty(const char(&name)[N]) const { char* d_ptr = detail::curve::DeviceCurve::getEnvironmentMacroProperty(name); #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS if (!d_ptr) { return DeviceMacroProperty(nullptr, nullptr); } return DeviceMacroProperty(reinterpret_cast(d_ptr), reinterpret_cast(d_ptr + (I * J * K * W * sizeof(T)))); // Read-write flag resides in 8 bits at the end of the buffer #else return DeviceMacroProperty(reinterpret_cast(d_ptr)); #endif } template __device__ __forceinline__ DeviceEnvironmentDirectedGraph ReadOnlyDeviceEnvironment::getDirectedGraph(const char(&name)[M]) const { const detail::curve::Curve::VariableHash graph_hash = detail::curve::Curve::variableHash(name); #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Seatbelts check that the graph exists, by looking for mandatory vertex _id property // Rather than it failing at property lookup time if (detail::curve::DeviceCurve::getVariableIndex((graph_hash ^ detail::curve::Curve::variableHash("_environment_directed_graph_vertex")) + detail::curve::Curve::variableHash("_id")) == detail::curve::DeviceCurve::UNKNOWN_VARIABLE) { DTHROW("Environment directed graph with name '%s' was not found\n", name); } #endif return DeviceEnvironmentDirectedGraph(graph_hash); } #else template __device__ __forceinline__ DeviceEnvironmentDirectedGraph ReadOnlyDeviceEnvironment::getDirectedGraph(const char(&name)[M]) const { const detail::curve::Curve::VariableHash graph_hash = detail::curve::Curve::variableHash(name); #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS // Seatbelts check that the graph exists, by requesting it's hash if (detail::curve::DeviceCurve::getGraphHash(name) == detail::curve::DeviceCurve::UNKNOWN_GRAPH) { DTHROW("Environment directed graph with name '%s' was not found\n", name); } #endif return DeviceEnvironmentDirectedGraph(graph_hash); } #endif // __CUDACC_RTC__ } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_DEVICEENVIRONMENT_CUH_