Packers

Packers are used to store data of a set of registered channels from a mirheo::DataManager into a single buffer and vice-versa. They are used to redistribute and exchange data accross neighbouring ranks efficiently. This allows to send single MPI messages instead of one message per channel.

Generic Packer

This is the base packer class. All packers contain generic packers that are used to pack different kind of data (such as particle or object data).

struct GenericPackerHandler

A device-friendly structure that is used to pack and unpack multiple channels into a single buffer.

Additionally to being packed and unpacked, the data can be shifted. This facilitate the exchange and redistribute operations.

The packed channels are structured in a single buffer containing:

1. The first channel data
2. padding
3. The second channel data
4. padding
5. …

Hence, the number of elements must be known in advance before packing. This is generally not a limitation, as memory must be allocated before packing.

Subclassed by mirheo::GenericPacker

Public Functions

size_t pack(int srcId, int dstId, char *dstBuffer, int numElements) const

Fetch one datum from the registered channels and pack it into a buffer.

Return

The size (in bytes) taken by the packed data (numElements elements)

Parameters

• srcId: Index of the datum to fetch from registered channel space (in number of elements).
• dstId: Index of the datum to store in dstBuffer space (in number of elements).
• dstBuffer: Destination buffer
• numElements: Total number of elements that will be packed in the buffer.

size_t packShift(int srcId, int dstId, char *dstBuffer, int numElements, real3 shift) const

Fetch one datum from the registered channels, shift it (if applicable) and pack it into the buffer.

Only channels with active shift will be shifted.

Return

The size (in bytes) taken by the packed data (numElements elements)

Parameters

• srcId: Index of the datum to fetch from registered channel space (in number of elements).
• dstId: Index of the datum to store in dstBuffer space (in number of elements).
• dstBuffer: Destination buffer
• numElements: Total number of elements that will be packed in the buffer.
• shift: The coordinate shift

size_t unpack(int srcId, int dstId, const char *srcBuffer, int numElements) const

Unpack one datum from the buffer and store it in the registered channels.

Return

The size (in bytes) taken by the packed data (numElements elements)

Parameters

• srcId: Index of the datum to fetch from the buffer (in number of elements).
• dstId: Index of the datum to store in the registered channels (in number of elements).
• srcBuffer: Source buffer that contains packed data.
• numElements: Total number of elements that are packed in the buffer.

size_t unpackAtomicAddNonZero(int srcId, int dstId, const char *srcBuffer, int numElements) const

Unpack one datum from the buffer and add it to the registered channels atomically.

Return

The size (in bytes) taken by the packed data (numElements elements)

Parameters

• srcId: Index of the datum to fetch from the buffer (in number of elements).
• dstId: Index of the datum to add to the registered channels (in number of elements).
• srcBuffer: Source buffer that contains packed data.
• numElements: Total number of elements that are packed in the buffer.

size_t unpackShift(int srcId, int dstId, const char *srcBuffer, int numElements, real3 shift) const

Unpack and shift one datum from the buffer and store it in the registered channels.

Return

The size (in bytes) taken by the packed data (numElements elements)

Parameters

• srcId: Index of the datum to fetch from the buffer (in number of elements).
• dstId: Index of the datum to store into the registered channels (in number of elements).
• srcBuffer: Source buffer that contains packed data.
• numElements: Total number of elements that are packed in the buffer.
• shift: The coordinate shift

void copyTo(GenericPackerHandler &dst, int srcId, int dstId) const

Copy one datum from the registered channels to the registered channels of another GenericPackerHandler.

Parameters

• dst: The other GenericPackerHandler that will receive the new datum.
• srcId: Index of the datum to fetch from the registered channels (in number of elements).
• dstId: Index of the datum to store into the dst registered channels (in number of elements).

size_t getSizeBytes(int numElements) const

Get the size (in bytes) of the buffer that can hold the packed data of numElements elements from all registered channels.

This must be used to allocate the buffer size. Because of padding, the size is not simply the sum of sizes of all elements.

Return

The size (in bytes) of the buffer.

Parameters

• numElements: The number of elements that the buffer must contain once packed.

Public Static Attributes

constexpr size_t alignment = getPaddedSize<char>(1)

Alignment sufficient for all types used in channels.

Useful for external codes that operate with Mirheo’s packing functions.

class GenericPacker : mirheo::GenericPackerHandler

This class is used to construct GenericPackerHandler, to be passed to the device.

Public Functions

void updateChannels(DataManager &dataManager, PackPredicate &predicate, cudaStream_t stream)

Register all channels of a DataManager satisfying a predicate.

All previously registered channels will be removed before adding those described above.

Parameters

• dataManager: The object that contains the channels to register
• predicate: The filter (white list) that is used to select the channels to register, based on their description and names
• stream: The stream used to transfer the data on the device

GenericPackerHandler &handler()

Get a handler that can be used on the device.

size_t getSizeBytes(int numElements) const

see GenericPackerHandler::getSizeBytes().

Particles Packer

struct ParticlePackerHandler

A packer specific to particle data only.

The user can use the internal generic packer directly.

Subclassed by mirheo::ObjectPackerHandler

Public Functions

size_t getSizeBytes(int numElements) const

Get the reuired size (in bytes) of the buffer to hold the packed data.

Public Members

GenericPackerHandler particles

The packer responsible for the particles.

class ParticlePacker

Helper class to construct a ParticlePackerHandler.

Subclassed by mirheo::ObjectPacker

Public Functions

ParticlePacker(PackPredicate predicate)

Construct a ParticlePacker.

Parameters

• predicate: The channel filter that will be used to select the channels to be registered.

virtual void update(LocalParticleVector *lpv, cudaStream_t stream)

Register the channels of a LocalParticleVector that meet the predicate requirements.

Parameters

• lpv: The LocalParticleVector that holds the channels to be registered.
• stream: The stream used to transfer the channels information to the device.

ParticlePackerHandler handler()

get a handler usable on device

virtual size_t getSizeBytes(int numElements) const

Get the reuired size (in bytes) of the buffer to hold all the packed data.

Objects Packer

struct ObjectPackerHandler : public mirheo::ParticlePackerHandler

A packer specific to objects.

Will store both particle and object data into a single buffer.

Subclassed by mirheo::RodPackerHandler

Public Functions

size_t getSizeBytes(int numElements) const

Get the reuired size (in bytes) of the buffer to hold the packed data.

__device__ size_t blockPack(int numElements, char *buffer, int srcObjId, int dstObjId) const

Fetch a full object from the registered channels and pack it into the buffer.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of objects that will be packed in the buffer.
• buffer: Destination buffer that will hold the packed object
• srcObjId: The index of the object to fetch from registered channels
• dstObjId: The index of the object to store into the buffer

__device__ size_t blockPackShift(int numElements, char *buffer, int srcObjId, int dstObjId, real3 shift) const

Fetch a full object from the registered channels, shift it and pack it into the buffer.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of objects that will be packed in the buffer.
• buffer: Destination buffer that will hold the packed object
• srcObjId: The index of the object to fetch from registered channels
• dstObjId: The index of the object to store into the buffer
• shift: The coordnate shift

__device__ size_t blockUnpack(int numElements, const char *buffer, int srcObjId, int dstObjId) const

Unpack a full object from the buffer and store it into the registered channels.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of objects that will be packed in the buffer.
• buffer: Buffer that holds the packed object
• srcObjId: The index of the object to fetch from the buffer
• dstObjId: The index of the object to store into the registered channels

__device__ size_t blockUnpackAddNonZero(int numElements, const char *buffer, int srcObjId, int dstObjId) const

Unpack a full object from the buffer and add it to the registered channels.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of objects that will be packed in the buffer.
• buffer: Buffer that holds the packed object
• srcObjId: The index of the object to fetch from the buffer
• dstObjId: The index of the object to store into the registered channels

__device__ size_t blockUnpackShift(int numElements, const char *buffer, int srcObjId, int dstObjId, real3 shift) const

Unpack a full object from the buffer, shift it and store it into the registered channels.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of objects that will be packed in the buffer.
• buffer: Buffer that holds the packed object
• srcObjId: The index of the object to fetch from the buffer
• dstObjId: The index of the object to store into the registered channels
• shift: Coordinates shift

__device__ void blockCopyParticlesTo(ParticlePackerHandler &dst, int srcObjId, int dstPartIdOffset) const

Copy the particle data of a full object from registered channels into the registered channels of a ParticlePackerHandler.

This method must be called by one CUDA block per object.

Parameters

• dst: The destination ParticlePackerHandler
• srcObjId: The index of the object to fetch from the registered channels
• dstPartIdOffset: The index of the first particle in the destination ParticlePackerHandler

__device__ void blockCopyTo(ObjectPackerHandler &dst, int srcObjId, int dstObjId) const

Copy a full object from registered channels into the registered channels of a ObjectPackerHandler.

This method must be called by one CUDA block per object.

Parameters

• dst: The destination ObjectPackerHandler
• srcObjId: The index of the object to fetch from the registered channels
• dstObjId: The index of the object to store in the destination ObjectPackerHandler

Public Members

int objSize

number of particles per object

GenericPackerHandler objects

packer responsible for the object data

class ObjectPacker : public mirheo::ParticlePacker

Helper class to construct a ObjectPackerHandler.

Subclassed by mirheo::RodPacker

Public Functions

ObjectPacker(PackPredicate predicate)

Construct a ObjectPacker.

Parameters

• predicate: The channel filter that will be used to select the channels to be registered.

void update(LocalParticleVector *lpv, cudaStream_t stream)

Register the channels of a LocalParticleVector that meet the predicate requirements.

Parameters

• lpv: The LocalParticleVector that holds the channels to be registered.
• stream: The stream used to transfer the channels information to the device.

ObjectPackerHandler handler()

get a handler usable on device

size_t getSizeBytes(int numElements) const

Get the reuired size (in bytes) of the buffer to hold all the packed data.

Rods Packer

struct RodPackerHandler : public mirheo::ObjectPackerHandler

A packer specific to rods.

Will store particle, object and bisegment data into a single buffer.

Public Functions

size_t getSizeBytes(int numElements) const

Get the reuired size (in bytes) of the buffer to hold the packed data.

__device__ size_t blockPack(int numElements, char *buffer, int srcObjId, int dstObjId) const

Fetch a full rod from the registered channels and pack it into the buffer.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements rods). Only relevant for thread with Id 0.

Parameters

• numElements: Number of rods that will be packed in the buffer.
• buffer: Destination buffer that will hold the packed rod
• srcObjId: The index of the rod to fetch from registered channels
• dstObjId: The index of the rod to store into the buffer

__device__ size_t blockPackShift(int numElements, char *buffer, int srcObjId, int dstObjId, real3 shift) const

Fetch a full rod from the registered channels, shift it and pack it into the buffer.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements rods). Only relevant for thread with Id 0.

Parameters

• numElements: Number of rods that will be packed in the buffer.
• buffer: Destination buffer that will hold the packed rod
• srcObjId: The index of the rod to fetch from registered channels
• dstObjId: The index of the rod to store into the buffer
• shift: The coordnate shift

__device__ size_t blockUnpack(int numElements, const char *buffer, int srcObjId, int dstObjId) const

Unpack a full rod from the buffer and store it into the registered channels.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of rods that will be packed in the buffer.
• buffer: Buffer that holds the packed rod
• srcObjId: The index of the rod to fetch from the buffer
• dstObjId: The index of the rod to store into the registered channels

__device__ size_t blockUnpackAddNonZero(int numElements, const char *buffer, int srcObjId, int dstObjId) const

Unpack a full rod from the buffer and add it into the registered channels.

This method must be called by one CUDA block per object.

Return

The size (in bytes) taken by the packed data (numElements objects). Only relevant for thread with Id 0.

Parameters

• numElements: Number of rods that will be packed in the buffer.
• buffer: Buffer that holds the packed rod
• srcObjId: The index of the rod to fetch from the buffer
• dstObjId: The index of the rod to store into the registered channels

Public Members

int nBisegments

number of bisegment per rod

GenericPackerHandler bisegments

packer responsible for the bisegment data

class RodPacker : public mirheo::ObjectPacker

Helper class to construct a RodPackerHandler.

Public Functions

RodPacker(PackPredicate predicate)

Construct a RodPacker.

Parameters

• predicate: The channel filter that will be used to select the channels to be registered.

void update(LocalParticleVector *lpv, cudaStream_t stream)

Register the channels of a LocalParticleVector that meet the predicate requirements.

Parameters

• lpv: The LocalParticleVector that holds the channels to be registered.
• stream: The stream used to transfer the channels information to the device.

RodPackerHandler handler()

get a handler usable on device

size_t getSizeBytes(int numElements) const

Get the reuired size (in bytes) of the buffer to hold all the packed data.