Class CUDAScatter

Class Documentation

class CUDAScatter

Singleton class for performing generic scatters This is used for optional messages, agent death, agent birth

Public Types

typedef CUDAScanCompaction::Type Type

Flag used to decide which scan_flag array should be used

See

CUDAScanCompaction::type

Public Functions

inline CUDAScanCompaction &Scan()
unsigned int scatter(const unsigned int &streamResourceId, const cudaStream_t &stream, const Type &messageOrAgent, const VariableMap &vars, const std::map<std::string, void*> &in, const std::map<std::string, void*> &out, const unsigned int &itemCount, const unsigned int &out_index_offset = 0, const bool &invert_scan_flag = false, const unsigned int &scatter_all_count = 0)

Convenience wrapper for scatter() Scatters agents from SoA to SoA according to d_position flag Used for device agent creation and agent death CUDAScanCompaction::scan_flag is used to decide who should be scattered CUDAScanCompaction::position is used to decide where to scatter to

Note

This is deprecated, unclear if still used

Parameters
  • streamResourceId – Index of internal resources to use

  • messageOrAgent – Flag of whether message or agent CUDAScanCompaction arrays should be used

  • varsVariable description map from ModelData hierarchy

  • in – Input variable name:ptr map @paramn out Output variable name:ptr map

  • itemCount – Total number of items in input array to consider

  • out_index_offset – The offset to be applied to the ouput index (e.g. if out already contains data) @parma invert_scan_flag If true, agents with scan_flag set to 0 will be moved instead @parma scatter_all_count The number of agents at the start of in to be copied, ones after this use scanflag

unsigned int scatter(const unsigned int &streamResourceId, const cudaStream_t &stream, const Type &messageOrAgent, const std::vector<ScatterData> &scatterData, const unsigned int &itemCount, const unsigned int &out_index_offset = 0, const bool &invert_scan_flag = false, const unsigned int &scatter_all_count = 0)

Scatters agents from SoA to SoA according to d_position flag Used for device agent creation and agent death CUDAScanCompaction::scan_flag is used to decide who should be scattered CUDAScanCompaction::position is used to decide where to scatter to

Parameters
  • streamResourceId – Index of internal resources to use

  • messageOrAgent – Flag of whether message or agent CUDAScanCompaction arrays should be used

  • scatterData – Vector of scatter configuration for each variable to be scattered

  • itemCount – Total number of items in input array to consider

  • out_index_offset – The offset to be applied to the ouput index (e.g. if out already contains data) @parma invert_scan_flag If true, agents with scan_flag set to 0 will be moved instead @parma scatter_all_count The number of agents at the start of in to be copied, ones after this use scanflag

void scatterPosition(const unsigned int &streamResourceId, const cudaStream_t &stream, const Type &messageOrAgent, const std::vector<ScatterData> &scatterData, const unsigned int &itemCount)

Scatters agents from SoA to SoA according to d_position flag as input_source, all variables are scattered Used for Host function sort agent CUDAScanCompaction::position is used to decide where to scatter to

Parameters
  • streamResourceId – Index of internal resources to use

  • messageOrAgent – Flag of whether message or agent CUDAScanCompaction arrays should be used

  • scatterData – Vector of scatter configuration for each variable to be scattered

  • itemCount – Total number of items in input array to consider

unsigned int scatterCount(const unsigned int &streamResourceId, const cudaStream_t &stream, const Type &messageOrAgent, const unsigned int &itemCount, const unsigned int &scatter_all_count = 0)

Returns the final CUDAScanCompaction::position item Same value as scatter, - scatter_a__count

Parameters
  • streamResourceId – Index of internal resources to use

  • messageOrAgent – Flag of whether message or agent CUDAScanCompaction arrays should be used

  • itemCount – Total number of items in input array to consider @parma scatter_all_count The number offset into the array where the scan began

unsigned int scatterAll(const unsigned int &streamResourceId, const cudaStream_t &stream, const std::vector<ScatterData> &scatterData, const unsigned int &itemCount, const unsigned int &out_index_offset = 0)

Scatters a contigous block from SoA to SoA CUDAScanCompaction::scan_flag/position are not used

Note

If calling scatter() with itemCount == scatter_all_count works the same

Parameters
  • streamResourceId – Index of internal resources to use

  • scatterData – Vector of scatter configuration for each variable to be scattered

  • itemCount – Total number of items in input array to consider

  • out_index_offset – The offset to be applied to the ouput index (e.g. if out already contains data)

unsigned int scatterAll(const unsigned int &streamResourceId, const cudaStream_t &stream, const VariableMap &vars, const std::map<std::string, void*> &in, const std::map<std::string, void*> &out, const unsigned int &itemCount, const unsigned int &out_index_offset)

Convenience wrapper to scatterAll()

Parameters
  • streamResourceId – Index of internal resources to use

  • varsVariable description map from ModelData hierarchy

  • in – Input variable name:ptr map @paramn out Output variable name:ptr map

void pbm_reorder(const unsigned int &streamResourceId, const cudaStream_t &stream, const VariableMap &vars, const std::map<std::string, void*> &in, const std::map<std::string, void*> &out, const unsigned int &itemCount, const unsigned int *d_bin_index, const unsigned int *d_bin_sub_index, const unsigned int *d_pbm)

Used for reordering messages from SoA to SoA Position information is taken using PBM data, rather than d_position Used by spatial messaging.

Parameters
  • streamResourceId – Index of internal resources to use

  • varsVariable description map from ModelData hierarchy

  • in – Input variable name:ptr map @paramn out Output variable name:ptr map

  • itemCount – Total number of items in input array to consider

  • d_bin_index – This idenitifies which bin each index should be sorted to

  • d_bin_sub_index – This indentifies where within it’s bin, an index should be sorted to

  • d_pbm – This is the PBM, it identifies at which index a bin’s storage begins

void scatterNewAgents(const unsigned int &streamResourceId, const cudaStream_t &stream, const std::vector<ScatterData> &scatterData, const size_t &totalAgentSize, const unsigned int &inCount, const unsigned int &out_index_offset)

Scatters agents from AoS to SoA Used by host agent creation

Parameters
  • streamResourceId – Index of internal resources to use

  • scatterData – Vector of scatter configuration for each variable to be scattered

  • totalAgentSize – Total size of all of the variables in an agent

  • inCount – Total number of items in input array to consider

  • out_index_offset – The offset to be applied to the ouput index (e.g. if out already contains data)

void broadcastInit(const unsigned int &streamResourceId, const cudaStream_t &stream, const std::list<std::shared_ptr<VariableBuffer>> &vars, const unsigned int &itemCount, const unsigned int out_index_offset)

Broadcasts a single value for each variable to a contiguous block in SoA Used prior to device agent creation

Parameters
  • streamResourceId – Index of internal resources to use

  • varsVariable description map from ModelData hierarchy

  • itemCount – Total number of items in input array to consider

  • out_index_offset – The offset to be applied to the ouput index (e.g. if out already contains data)

void broadcastInit(const unsigned int &streamResourceId, const cudaStream_t &stream, const VariableMap &vars, void *const d_newBuff, const unsigned int &itemCount, const unsigned int out_index_offset)
void arrayMessageReorder(const unsigned int &streamResourceId, const cudaStream_t &stream, const VariableMap &vars, const std::map<std::string, void*> &in, const std::map<std::string, void*> &out, const unsigned int &itemCount, const unsigned int &array_length, unsigned int *d_write_flag = nullptr)

Used to reorder array messages based on __INDEX variable, that variable is not sorted Also throws exception if any indexes are repeated

Parameters
  • streamResourceId – Index of internal resources to use

  • array_length – Length of the array messages are to be stored in (max index + 1)

  • d_write_flag – Device pointer to array for tracking how many messages output to each bin, caller responsibiltiy to ensure it is array_length or longer

inline CUDAScatter()
void purge()

Wipes out host mirrors of device memory Only really to be used after calls to cudaDeviceReset()

Note

Only currently used after some tests

CUDAScatter(CUDAScatter const&) = delete
void operator=(CUDAScatter const&) = delete
struct InversionIterator

This utility class provides a wrapper for unsigned int * When the iterator is dereferenced the pointed to unsigned int is evaluated using invert() This is useful when trying to partition and sort a dataset using only scatter and scan

Public Types

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

Public Functions

__host__ __device__ inline explicit InversionIterator(unsigned int *_p)
__device__ inline InversionIterator &operator=(const InversionIterator &other)
__device__ inline InversionIterator operator++(int a)
__device__ inline InversionIterator operator++()
__device__ inline unsigned int operator*()
__device__ inline InversionIterator operator+(const int &b) const
__device__ inline unsigned int operator[](int b) const
struct ScatterData

As we scatter per variable, this structure holds all the data required for a single variable

Public Members

size_t typeLen
char *const in
char *out