Program Listing for File HostMacroProperty.cuh

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

#ifndef INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_HOSTMACROPROPERTY_CUH_
#define INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_HOSTMACROPROPERTY_CUH_

#include <cstdint>
#include <algorithm>
#include <array>
#include <memory>
#include <cmath>
#include <string>

namespace flamegpu {

struct HostMacroProperty_MetaData {
    HostMacroProperty_MetaData(void* _d_base_ptr, const std::array<unsigned int, 4>& _dims, size_t _type_size,
        bool _device_read_flag, const std::string &name, cudaStream_t _stream)
        : h_base_ptr(nullptr)
        , d_base_ptr(static_cast<char*>(_d_base_ptr))
        , dims(_dims)
        , elements(dims[0] * dims[1] * dims[2] * dims[3])
        , type_size(_type_size)
        , has_changed(false)
        , device_read_flag(_device_read_flag)
        , property_name(name)
        , stream(_stream)
    { }
    ~HostMacroProperty_MetaData() {
        upload();
        if (h_base_ptr)
            std::free(h_base_ptr);
    }
    void download() {
        if (!h_base_ptr) {
            force_download();
        }
    }

    void force_download() {
        if (!h_base_ptr) {
            h_base_ptr = static_cast<char*>(malloc(elements * type_size));
        }
        gpuErrchk(cudaMemcpyAsync(h_base_ptr, d_base_ptr, elements * type_size, cudaMemcpyDeviceToHost, stream));
        gpuErrchk(cudaStreamSynchronize(stream));
        has_changed = false;
    }
    void upload() {
        if (h_base_ptr && has_changed) {
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
            if (device_read_flag) {
                THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, "
                    "in HostMacroProperty_MetaData::upload()\n",
                    property_name.c_str());
            }
#endif
            gpuErrchk(cudaMemcpyAsync(d_base_ptr, h_base_ptr, elements * type_size, cudaMemcpyHostToDevice, stream));
            gpuErrchk(cudaStreamSynchronize(stream));
            has_changed = false;
        }
    }
    char* h_base_ptr;
    char* d_base_ptr;
    std::array<unsigned int, 4> dims;
    unsigned int elements;
    size_t type_size;
    bool has_changed;
    bool device_read_flag;
    std::string property_name;
    cudaStream_t stream;
};

template<typename T, unsigned int I = 1, unsigned int J = 1, unsigned int K = 1, unsigned int W = 1>
class HostMacroProperty {
    ptrdiff_t offset;
    std::shared_ptr<HostMacroProperty_MetaData> metadata;

 public:
    explicit HostMacroProperty(const std::shared_ptr<HostMacroProperty_MetaData>& _metadata);
    HostMacroProperty(const std::shared_ptr<HostMacroProperty_MetaData> &_metadata, const ptrdiff_t& _offset);

    HostMacroProperty<T, J, K, W, 1> operator[](unsigned int i) const;
    operator T();
    operator T() const;
    HostMacroProperty<T, I, J, K, W>& operator=(T val);
    void zero();

    HostMacroProperty<T, I, J, K, W>& operator++();
    T operator++(int);
    HostMacroProperty<T, I, J, K, W>& operator--();
    T operator--(int);

// Swig breaks if it see's auto return type
#ifndef SWIG
    template<typename T2>
    auto operator+(T2 b) const;
    template<typename T2>
    auto operator-(T2 b) const;
    template<typename T2>
    auto operator*(T2 b) const;
    template<typename T2>
    auto operator/(T2 b) const;
    template<typename T2>
    auto operator%(T2 b) const;
#endif
    // HostMacroProperty<T, I, J, K, W>& operator=(T b);  // Defined above
    template<typename T2>
    HostMacroProperty<T, I, J, K, W>& operator+=(T2 b);
    template<typename T2>
    HostMacroProperty<T, I, J, K, W>& operator-=(T2 b);
    template<typename T2>
    HostMacroProperty<T, I, J, K, W>& operator*=(T2 b);
    template<typename T2>
    HostMacroProperty<T, I, J, K, W>& operator/=(T2 b);
    template<typename T2>
    HostMacroProperty<T, I, J, K, W>& operator%=(T2 b);

 private:
    T& _get();
    T& _get() const;
};

#ifdef SWIG
template<typename T>
class HostMacroProperty_swig {
    const std::array<unsigned int, 4> dimensions;
    ptrdiff_t offset;
    std::shared_ptr<HostMacroProperty_MetaData> metadata;

 public:
    explicit HostMacroProperty_swig(const std::shared_ptr<HostMacroProperty_MetaData> &_metadata);
    HostMacroProperty_swig(const std::shared_ptr<HostMacroProperty_MetaData>& _metadata, const ptrdiff_t& _offset, const std::array<unsigned int, 4>& _dimensions);
    HostMacroProperty_swig<T> __getitem__(unsigned int i) const;
    void __setitem__(unsigned int i,  T val);
    void set(T val);

    int __int__();
    int64_t __long__();
    double __float__();
    bool __bool__();

    bool __eq__(T other) const;
    bool __ne__(T other) const;
    bool __lt__(T other) const;
    bool __le__(T other) const;
    bool __gt__(T other) const;
    bool __ge__(T other) const;

    unsigned int __len__() const { return dimensions[0]; }

    T get() const;
    void zero();
};
#endif

template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>::HostMacroProperty(const std::shared_ptr<HostMacroProperty_MetaData> &_metadata)
: offset(0u)
, metadata(_metadata)
{ }
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>::HostMacroProperty(const std::shared_ptr<HostMacroProperty_MetaData> &_metadata, const ptrdiff_t & _offset)
    : offset(_offset)
    , metadata(_metadata)
{ }
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, J, K, W, 1> HostMacroProperty<T, I, J, K, W>::operator[](unsigned int i) const {
    if (I == 1 && J == 1 && K == 1 && W == 1) {
        THROW exception::InvalidOperation("Indexing error, property has less dimensions.\n");
    } else if (i >= I) {
        THROW exception::OutOfBoundsException("Indexing error, out of bounds %u >= %u.\n", i, I);
    }
    // (i * J * K * W) + (j * K * W) + (k * W) + w
    return HostMacroProperty<T, J, K, W, 1>(metadata, offset + (i * J * K * W));
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>::operator T() {
    if (I != 1 || J != 1 || K != 1 || W != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    // Must assume changed
    metadata->has_changed = true;
    return *(reinterpret_cast<T*>(metadata->h_base_ptr) + offset);
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>::operator T() const {
    metadata->download();
    return *(reinterpret_cast<T*>(metadata->h_base_ptr) + offset);
}

template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
void HostMacroProperty<T, I, J, K, W>::zero() {
    if (metadata->h_base_ptr) {
        // Memset on host
        memset(reinterpret_cast<T*>(metadata->h_base_ptr) + offset, 0, I * J * K * W * metadata->type_size);
        metadata->has_changed = true;
    } else {
#if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
        if (metadata->device_read_flag) {
            THROW flamegpu::exception::InvalidEnvProperty("The environment macro property '%s' was not found, "
                "in HostMacroProperty::zero()\n",
                metadata->property_name.c_str());
        }
#endif
        // Memset on device
        gpuErrchk(cudaMemset(reinterpret_cast<T*>(metadata->d_base_ptr) + offset, 0, I * J * K * W * metadata->type_size));
    }
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
T& HostMacroProperty<T, I, J, K, W>::_get() const {
    if (I != 1 || J != 1 || K != 1 || W != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset];
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
T& HostMacroProperty<T, I, J, K, W>::_get() {
    if (I != 1 || J != 1 || K != 1 || W != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    metadata->has_changed = true;
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset];
}


template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator++() {
    ++_get();
    return *this;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
T HostMacroProperty<T, I, J, K, W>::operator++(int) {
    T &t = _get();
    T ret = t;
    ++t;
    return ret;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator--() {
    --_get();
    return *this;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
T HostMacroProperty<T, I, J, K, W>::operator--(int) {
    T& t = _get();
    T ret = t;
    --t;
    return ret;
}

#ifndef SWIG
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
auto HostMacroProperty<T, I, J, K, W>::operator+(const T2 b) const {
    return _get() + b;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
auto HostMacroProperty<T, I, J, K, W>::operator-(const T2 b) const {
    return _get() - b;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
auto HostMacroProperty<T, I, J, K, W>::operator*(const T2 b) const {
    return _get() * b;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
auto HostMacroProperty<T, I, J, K, W>::operator/(const T2 b) const {
    return _get() / b;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
auto HostMacroProperty<T, I, J, K, W>::operator%(const T2 b) const {
    return _get() % b;
}
#endif

template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator+=(const T2 b) {
    _get() += b;
    return *this;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator-=(const T2 b) {
    _get() -= b;
    return *this;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator*=(const T2 b) {
    _get() *= b;
    return *this;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator/=(const T2 b) {
    _get() /= b;
    return *this;
}
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
template<typename T2>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator%=(const T2 b) {
    _get() %= b;
    return *this;
}

template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W>
HostMacroProperty<T, I, J, K, W>& HostMacroProperty<T, I, J, K, W>::operator=(const T val) {
    if (I != 1 || J != 1 || K != 1 || W != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    reinterpret_cast<T*>(metadata->h_base_ptr)[offset] = val;
    metadata->has_changed = true;
    return *this;
}

// SWIG versions
#ifdef SWIG
template<typename T>
HostMacroProperty_swig<T>::HostMacroProperty_swig(const std::shared_ptr<HostMacroProperty_MetaData>& _metadata)
    : dimensions(_metadata->dims)
    , offset(0u)
    , metadata(_metadata)
{ }
template<typename T>
HostMacroProperty_swig<T>::HostMacroProperty_swig(const std::shared_ptr<HostMacroProperty_MetaData>& _metadata, const ptrdiff_t& _offset, const std::array<unsigned int, 4> &_dimensions)
    : dimensions(_dimensions)
    , offset(_offset)
    , metadata(_metadata)
{ }
template<typename T>
HostMacroProperty_swig<T> HostMacroProperty_swig<T>::__getitem__(unsigned int i) const {
    if (dimensions[0] == 1 && dimensions[1] == 1 && dimensions[2] == 1 && dimensions[3] == 1) {
        THROW exception::InvalidOperation("Indexing error, property has less dimensions.\n");
    } else if (i >= dimensions[0]) {
        THROW exception::OutOfBoundsException("Indexing error, out of bounds %u >= %u.\n", i, dimensions[0]);
    }
    // (i * J * K * W) + (j * K * W) + (k * W) + w
    return HostMacroProperty_swig<T>(metadata, offset + (i * dimensions[1] * dimensions[2] * dimensions[3]), { dimensions[1], dimensions[2], dimensions[3], 1 });
}

template<typename T>
void HostMacroProperty_swig<T>::zero() {
    if (metadata->h_base_ptr) {
        // Memset on host
        memset(reinterpret_cast<T*>(metadata->h_base_ptr) + offset, 0, dimensions[0] * dimensions[1] * dimensions[2] * dimensions[3] * metadata->type_size);
        metadata->has_changed = true;
    } else {
        // Memset on device
        gpuErrchk(cudaMemset(reinterpret_cast<T*>(metadata->d_base_ptr) + offset, 0, dimensions[0] * dimensions[1] * dimensions[2] * dimensions[3] * metadata->type_size));
    }
}
template<typename T>
void HostMacroProperty_swig<T>::set(T val) {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    reinterpret_cast<T*>(metadata->h_base_ptr)[offset] = val;
    metadata->has_changed = true;
}
template<typename T>
void HostMacroProperty_swig<T>::__setitem__(unsigned int i, const T val) {
    if (dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    } else if (i >= dimensions[0]) {
        THROW exception::InvalidOperation("Indexing out of bounds %u >= %u.\n", i, dimensions[0]);
    }
    metadata->download();
    unsigned int t_offset = offset + (i * dimensions[1] * dimensions[2] * dimensions[3]);
    reinterpret_cast<T*>(metadata->h_base_ptr)[t_offset] = val;
    metadata->has_changed = true;
}
template<typename T>
int HostMacroProperty_swig<T>::__int__() {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return static_cast<int>(reinterpret_cast<T*>(metadata->h_base_ptr)[offset]);
}
template<typename T>
int64_t HostMacroProperty_swig<T>::__long__() {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return static_cast<int64_t>(reinterpret_cast<T*>(metadata->h_base_ptr)[offset]);
}
template<typename T>
double HostMacroProperty_swig<T>::__float__() {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return static_cast<double>(reinterpret_cast<T*>(metadata->h_base_ptr)[offset]);
}
template<typename T>
bool HostMacroProperty_swig<T>::__bool__() {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return static_cast<bool>(reinterpret_cast<T*>(metadata->h_base_ptr)[offset]);
}
template<typename T>
bool HostMacroProperty_swig<T>::__eq__(const T other) const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] == other;
}
template<typename T>
bool HostMacroProperty_swig<T>::__ne__(const T other) const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] != other;
}
template<typename T>
bool HostMacroProperty_swig<T>::__lt__(const T other) const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] < other;
}
template<typename T>
bool HostMacroProperty_swig<T>::__le__(const T other) const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] <= other;
}
template<typename T>
bool HostMacroProperty_swig<T>::__gt__(const T other) const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] > other;
}
template<typename T>
bool HostMacroProperty_swig<T>::__ge__(const T other) const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] >= other;
}
// template<typename T>
// T HostMacroProperty_swig<T>::__mod__(const T other) {
//     if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
//         THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
//     }
//     metadata->download();
//     return reinterpret_cast<T*>(metadata->h_base_ptr)[offset] % other;
// }
template<typename T>
T HostMacroProperty_swig<T>::get() const {
    if (dimensions[0] != 1 || dimensions[1] != 1 || dimensions[2] != 1 || dimensions[3] != 1) {
        THROW exception::InvalidOperation("Indexing error, property has more dimensions.\n");
    }
    metadata->download();
    return reinterpret_cast<T*>(metadata->h_base_ptr)[offset];
}
#endif

}  // namespace flamegpu

#endif  // INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_HOSTMACROPROPERTY_CUH_