.. _program_listing_file_include_flamegpu_simulation_detail_CUDAScatter.cuh:

Program Listing for File CUDAScatter.cuh
========================================

|exhale_lsh| :ref:`Return to documentation for file <file_include_flamegpu_simulation_detail_CUDAScatter.cuh>` (``include/flamegpu/simulation/detail/CUDAScatter.cuh``)

.. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS

.. code-block:: cpp

   #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_