Program Listing for File CUDAScatter.cuh

↰ Return to documentation for file (include/flamegpu/simulation/detail/CUDAScatter.cuh)

#ifndef INCLUDE_FLAMEGPU_SIMULATION_DETAIL_CUDASCATTER_CUH_
#define INCLUDE_FLAMEGPU_SIMULATION_DETAIL_CUDASCATTER_CUH_

#include <map>
#include <string>
#include <array>
#include <memory>
#include <list>
#include <vector>

#include "flamegpu/model/Variable.h"
#include "flamegpu/simulation/detail/CubTemporaryMemory.cuh"
#include "flamegpu/simulation/detail/CUDAScanCompaction.h"

namespace flamegpu {
struct VarOffsetStruct;
namespace detail {
struct VariableBuffer;

class CUDAScatter {
 public:
     struct InversionIterator {
         using difference_type = unsigned int;
         using value_type = unsigned int;
         using pointer = unsigned int*;
         using reference = unsigned int&;
         using iterator_category = std::random_access_iterator_tag;
         __host__ __device__ explicit InversionIterator(unsigned int *_p) : p(_p) { }

         __device__ InversionIterator &operator=(const InversionIterator&other) { p = other.p; return *this; }
         __device__ InversionIterator operator++ (int a) { p += a;  return *this; }
         __device__ InversionIterator operator++ () { p++;  return *this; }
         __device__ unsigned int operator *() { return invert(*p); }
         __device__ InversionIterator operator+(const int b) const { return InversionIterator(p + b); }
         __device__ unsigned int operator[](int b) const { return  invert(p[b]); }
      private:
         __device__ unsigned int invert(unsigned int c) const { return c == 0 ? 1 : 0; }
         unsigned int *p;
     };
    typedef CUDAScanCompaction::Type Type;
    struct ScatterData {
        size_t typeLen;
        char *const in;
        char *out;
    };

 private:
    struct StreamData {
        friend class std::array<StreamData, CUDAScanCompaction::MAX_STREAMS>;
        ScatterData *d_data;
        unsigned int data_len;
        StreamData();
        ~StreamData();
        void resize(unsigned int newLen);
    };
    std::array<StreamData, CUDAScanCompaction::MAX_STREAMS> streamResources;
    std::array<detail::CubTemporaryMemory, CUDAScanCompaction::MAX_STREAMS> cubTemps;

 public:
    CUDAScanCompaction &Scan() { return scan; }
    detail::CubTemporaryMemory &CubTemp(const unsigned int streamId) { return cubTemps[streamId]; }
     unsigned int scatter(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const Type &messageOrAgent,
        const VariableMap &vars,
        const std::map<std::string, void*> &in,
        const std::map<std::string, void*> &out,
        unsigned int itemCount,
        unsigned int out_index_offset = 0,
        bool invert_scan_flag = false,
        unsigned int scatter_all_count = 0);
    unsigned int scatter(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const Type &messageOrAgent,
        const std::vector<ScatterData> &scatterData,
        unsigned int itemCount,
        unsigned int out_index_offset = 0,
        bool invert_scan_flag = false,
        unsigned int scatter_all_count = 0);
    void scatterPosition_async(
        unsigned int streamResourceId,
        cudaStream_t stream,
        Type messageOrAgent,
        const std::vector<ScatterData>& scatterData,
        unsigned int itemCount);
    void scatterPosition_async(
        unsigned int streamResourceId,
        cudaStream_t stream,
        unsigned int *position,
        const std::vector<ScatterData>& scatterData,
        unsigned int itemCount);
    void scatterPosition(
        unsigned int streamResourceId,
        cudaStream_t stream,
        Type messageOrAgent,
        const std::vector<ScatterData> &scatterData,
        unsigned int itemCount);
    unsigned int scatterCount(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const Type &messageOrAgent,
        unsigned int itemCount,
        unsigned int scatter_all_count = 0);
    unsigned int scatterAll(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const std::vector<ScatterData> &scatterData,
        unsigned int itemCount,
        unsigned int out_index_offset = 0);
    unsigned int scatterAll(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const VariableMap &vars,
        const std::map<std::string, void*> &in,
        const std::map<std::string, void*> &out,
        unsigned int itemCount,
        unsigned int out_index_offset);
    void pbm_reorder(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const VariableMap &vars,
        const std::map<std::string, void*> &in,
        const std::map<std::string, void*> &out,
        unsigned int itemCount,
        const unsigned int *d_bin_index,
        const unsigned int *d_bin_sub_index,
        const unsigned int *d_pbm);
    void scatterNewAgents(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const std::vector<ScatterData> &scatterData,
        size_t totalAgentSize,
        unsigned int inCount,
        unsigned int out_index_offset);
    void broadcastInit_async(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const std::list<std::shared_ptr<VariableBuffer>> &vars,
        unsigned int itemCount,
        unsigned int out_index_offset);
    void broadcastInit(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const std::list<std::shared_ptr<VariableBuffer>>& vars,
        unsigned int itemCount,
        unsigned int out_index_offset);
    void broadcastInit_async(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const VariableMap& vars,
        void* const d_newBuff,
        unsigned int itemCount,
        unsigned int out_index_offset);
    void broadcastInit(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const VariableMap &vars,
        void * const d_newBuff,
        unsigned int itemCount,
        unsigned int out_index_offset);
    void arrayMessageReorder(
        unsigned int streamResourceId,
        cudaStream_t stream,
        const VariableMap &vars,
        const std::map<std::string, void*> &in,
        const std::map<std::string, void*> &out,
        unsigned int itemCount,
        unsigned int array_length,
        unsigned int *d_write_flag = nullptr);

 private:
    CUDAScanCompaction scan;

 public:
    CUDAScatter() { }
    // Public deleted creates better compiler errors
    CUDAScatter(CUDAScatter const&) = delete;
    void operator=(CUDAScatter const&) = delete;
};

}  // namespace detail
}  // namespace flamegpu

#endif  // INCLUDE_FLAMEGPU_SIMULATION_DETAIL_CUDASCATTER_CUH_