Program Listing for File DeviceEnvironment.cuh

↰ Return to documentation for file (include/flamegpu/runtime/environment/DeviceEnvironment.cuh)

#ifndef INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_DEVICEENVIRONMENT_CUH_
#define INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_DEVICEENVIRONMENT_CUH_

// #include <cuda_runtime.h>
#include <string>
#include <cassert>

#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<typename T, unsigned int M>
    __device__ __forceinline__ T getProperty(const char(&name)[M]) const;
    template<typename T, unsigned int N = 0, unsigned int M>
    __device__ __forceinline__ T getProperty(const char(&name)[M], unsigned int index) const;
    template<typename T, unsigned int I = 1, unsigned int J = 1, unsigned int K = 1, unsigned int W = 1, unsigned int M>
    __device__ __forceinline__ ReadOnlyDeviceMacroProperty<T, I, J, K, W> getMacroProperty(const char(&name)[M]) const;
    template<unsigned int M>
    __device__ __forceinline__ DeviceEnvironmentDirectedGraph getDirectedGraph(const char(&name)[M]) const;
};
class DeviceEnvironment : public ReadOnlyDeviceEnvironment {
 public:
    template<typename T, unsigned int I = 1, unsigned int J = 1, unsigned int K = 1, unsigned int W = 1, unsigned int M>
    __device__ __forceinline__ DeviceMacroProperty<T, I, J, K, W> 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<typename T, unsigned int M>
__device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[M]) const {
    return detail::curve::DeviceCurve::getEnvironmentProperty<T>(name);
}
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T ReadOnlyDeviceEnvironment::getProperty(const char(&name)[M], const unsigned int index) const {
    return detail::curve::DeviceCurve::getEnvironmentArrayProperty<T, N>(name,  index);
}

template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W, unsigned int N>
__device__ __forceinline__ ReadOnlyDeviceMacroProperty<T, I, J, K, W> ReadOnlyDeviceEnvironment::getMacroProperty(const char(&name)[N]) const {
    char * d_ptr = detail::curve::DeviceCurve::getEnvironmentMacroProperty<T, I, J, K, W>(name);
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (!d_ptr) {
        return ReadOnlyDeviceMacroProperty<T, I, J, K, W>(nullptr, nullptr);
    }
    return ReadOnlyDeviceMacroProperty<T, I, J, K, W>(reinterpret_cast<T*>(d_ptr),
        reinterpret_cast<unsigned int*>(d_ptr + (I * J * K * W * sizeof(T))));  // Read-write flag resides in 8 bits at the end of the buffer
#else

    return ReadOnlyDeviceMacroProperty<T, I, J, K, W>(reinterpret_cast<T*>(d_ptr));
#endif
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W, unsigned int N>
__device__ __forceinline__ DeviceMacroProperty<T, I, J, K, W> DeviceEnvironment::getMacroProperty(const char(&name)[N]) const {
    char* d_ptr = detail::curve::DeviceCurve::getEnvironmentMacroProperty<T, I, J, K, W>(name);
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
    if (!d_ptr) {
        return DeviceMacroProperty<T, I, J, K, W>(nullptr, nullptr);
    }
    return DeviceMacroProperty<T, I, J, K, W>(reinterpret_cast<T*>(d_ptr),
        reinterpret_cast<unsigned int*>(d_ptr + (I * J * K * W * sizeof(T))));  // Read-write flag resides in 8 bits at the end of the buffer
#else
    return DeviceMacroProperty<T, I, J, K, W>(reinterpret_cast<T*>(d_ptr));
#endif
}
template<unsigned int M>
__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<unsigned int M>
__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_