CollectiveDeviceBuffer< T > Class Template Reference

#include <CollectiveDeviceBuffer.h>

Public Member Functions
void	allocate (CollectiveBufferType type_in, const size_t numElemsIn, SynchronousCollectiveScope scope=SynchronousCollectiveScope::all)
	Allocates a symmetric buffer on all devices. More...
void	allocate_no_check (CollectiveBufferType type_in, const size_t numElemsIn)
	Allocates a symmetric buffer on all devices without reducing the buffer sizes. More...
void	reallocate (CollectiveBufferType type_in, const size_t newNumElems, const double factor, SynchronousCollectiveScope scope=SynchronousCollectiveScope::all)
	Reallocates a symmetric device buffer on all devices if needed. More...
void	reallocate_no_check (CollectiveBufferType type_in, const size_t newNumElems, const double factor)
void	deallocate ()
T *const	getDevicePtr () const
	Returns the pointer to the device buffer. More...
T **const	getDevicePeerPtr () const
	Returns the pointer to peer's pointer on the device. More...
std::vector< T * > const &	getHostPeer () const
	Returns a host-vector containing peer's device pointers. More...
CollectiveDeviceBuffer< T > &	operator= (CollectiveDeviceBuffer< T > &&other)
CollectiveDeviceBuffer< T > &	operator= (const CollectiveDeviceBuffer< T > &other)
	CollectiveDeviceBuffer (const CollectiveDeviceBuffer &other)
	CollectiveDeviceBuffer ()

Detailed Description

template<typename T>
class CollectiveDeviceBuffer< T >

Definition at line 25 of file CollectiveDeviceBuffer.h.

Constructor & Destructor Documentation

CollectiveDeviceBuffer() [1/2]

template<typename T>

CollectiveDeviceBuffer< T >::CollectiveDeviceBuffer ( const CollectiveDeviceBuffer< T > &  other )

inline

Definition at line 117 of file CollectiveDeviceBuffer.h.

                                                              {
    type = other.type;
    numElemsAlloc = other.numElemsAlloc;
    buffer = other.buffer;
    d_peerBuffers = other.d_peerBuffers;
    h_peerBuffers = other.h_peerBuffers;
  }

CollectiveDeviceBuffer() [2/2]

template<typename T>

CollectiveDeviceBuffer< T >::CollectiveDeviceBuffer ( )

inline

Definition at line 126 of file CollectiveDeviceBuffer.h.

{}

Member Function Documentation

allocate()

template<typename T >

void CollectiveDeviceBuffer< T >::allocate	(	CollectiveBufferType	type_in,
		const size_t	numElemsIn,
		SynchronousCollectiveScope	scope = SynchronousCollectiveScope::all
	)

Allocates a symmetric buffer on all devices.

This will allocate a buffer on all devices using the allocation method of the given CollectiveBufferType. The buffer will have the same size across all devices and this function will automatically compute the maximum number of elements across all master PEs.

This function can be called by either all PEs or just master PEs depending on the given scope

Definition at line 10 of file CollectiveDeviceBuffer.C.

References SynchronousCollectives::allReduce(), GlobalGPUMgr::getIsMasterPe(), GlobalGPUMgr::Object(), and SynchronousCollectives::Object().

                                    {
  SynchronousCollectives* syncColl = SynchronousCollectives::Object();
  GlobalGPUMgr* globalGPUMgr = GlobalGPUMgr::Object();
  const bool isMasterPe = globalGPUMgr->getIsMasterPe();

  size_t numElemsTemp;
  if (isMasterPe) { 
    numElemsTemp = numElemsIn;
  } else {
    numElemsTemp = 0;
  }

  // Compute size of buffer
  std::vector<size_t> numElemsVec = {numElemsTemp};
  const size_t maxNumElems = syncColl->allReduce<size_t>(numElemsVec, CkReduction::max_ulong_long,
    scope)[0];

  // Call the no check version of allocate
  allocate_no_check(type_in, maxNumElems);
}

allocate_no_check()

template<typename T >

void CollectiveDeviceBuffer< T >::allocate_no_check	(	CollectiveBufferType	type_in,
		const size_t	numElemsIn
	)

Allocates a symmetric buffer on all devices without reducing the buffer sizes.

This function will allocate a buffer on all devices, but it assumes that the number of elements given to it is the same across all master pes

Definition at line 33 of file CollectiveDeviceBuffer.C.

References SynchronousCollectives::allGather(), cudaCheck, GlobalGPUMgr::getDeviceIndex(), GlobalGPUMgr::getIsMasterPe(), GlobalGPUMgr::getNumDevices(), IPC, master, NAMD_die(), GlobalGPUMgr::Object(), SynchronousCollectives::Object(), and SingleProcess.

                                                                                                       {
  SynchronousCollectives* syncColl = SynchronousCollectives::Object();
  GlobalGPUMgr* globalGPUMgr = GlobalGPUMgr::Object();
  const bool isMasterPe = globalGPUMgr->getIsMasterPe();
  const int numDevices = globalGPUMgr->getNumDevices();
  const int deviceIndex = globalGPUMgr->getDeviceIndex();

  type = type_in;

  numElemsAlloc = numElemsIn;
  
  if (CollectiveBufferType::SingleProcess == type) {
    if (isMasterPe) {
      // Allocate Buffer locally
      allocate_device<T>(&(buffer), numElemsAlloc);

      // Communicate buffer to peers
      auto temp_peerBuffers = syncColl->allGather<unsigned long long>(
          (unsigned long long) buffer, SynchronousCollectiveScope::master);

      // Copy peer buffers to vector
      h_peerBuffers.resize(numDevices);
      for (int i = 0; i < numDevices; i++) {
        h_peerBuffers[i] = (T*) temp_peerBuffers[i];
      }

      // Copy P2P buffer to device
      allocate_device<T*>(&(d_peerBuffers), numDevices);
      copy_HtoD<T*>(h_peerBuffers.data(), d_peerBuffers, numDevices, nullptr);
      cudaStreamSynchronize(nullptr);
    }
#if ! defined(NAMD_HIP)
  } else if (CollectiveBufferType::IPC == type) {
    if (isMasterPe) {
      cudaIpcMemHandle_t handle;
      // Allocate Buffer locally
      allocate_device<T>(&(buffer), numElemsAlloc);
      cudaCheck(cudaIpcGetMemHandle(&handle, (void*) buffer));

      // Communicate handles to peers
      auto temp_peerHandles = syncColl->allGather<cudaIpcMemHandle_t>(handle, 
        SynchronousCollectiveScope::master);

      // Open CUDA IPC mem handles and store in vector
      h_peerBuffers.resize(numDevices);
      for (int i = 0; i < numDevices; i++) {
        if (i != deviceIndex) {
          cudaCheck(cudaIpcOpenMemHandle((void**) &(h_peerBuffers[i]), 
            temp_peerHandles[i],
            cudaIpcMemLazyEnablePeerAccess));
        } else {
          h_peerBuffers[i] = buffer;
        }
      }

      // Copy P2P buffer to device
      allocate_device<T*>(&(d_peerBuffers), numDevices);
      copy_HtoD<T*>(h_peerBuffers.data(), d_peerBuffers, numDevices, nullptr);
      cudaCheck(cudaStreamSynchronize(nullptr));
    }
#endif
  } else {
    NAMD_die("CollectiveBufferType not currently implemented");
  }
}

deallocate()

template<typename T >

void CollectiveDeviceBuffer< T >::deallocate

(

)

Definition at line 145 of file CollectiveDeviceBuffer.C.

References GlobalGPUMgr::getIsMasterPe(), and GlobalGPUMgr::Object().

                                           {
  // This is needed because different PEs can call functions on objects created by other PEs.
  // When this happens, both PEs could try to call deallocate leading to a double free and
  // a seg fault. This seems like a hacky way to fix
  GlobalGPUMgr* globalGPUMgr = GlobalGPUMgr::Object();

  numElemsAlloc = 0;

  if (!globalGPUMgr->getIsMasterPe()) return;

  if (buffer) {
    deallocate_device<T>(&buffer);
    buffer = nullptr;
  }
  if (d_peerBuffers) {
    deallocate_device<T*>(&d_peerBuffers);
    d_peerBuffers = nullptr;
  }
}

getDevicePeerPtr()

template<typename T>

T** const CollectiveDeviceBuffer< T >::getDevicePeerPtr ( ) const

inline

Returns the pointer to peer's pointer on the device.

Definition at line 76 of file CollectiveDeviceBuffer.h.

                                     {
    return d_peerBuffers;
  }

getDevicePtr()

template<typename T>

T* const CollectiveDeviceBuffer< T >::getDevicePtr ( ) const

inline

Returns the pointer to the device buffer.

Definition at line 69 of file CollectiveDeviceBuffer.h.

                                 {
    return buffer;
  }

getHostPeer()

template<typename T>

std::vector<T*> const& CollectiveDeviceBuffer< T >::getHostPeer ( ) const

inline

Returns a host-vector containing peer's device pointers.

Definition at line 83 of file CollectiveDeviceBuffer.h.

                                           {
    return h_peerBuffers;
  }

operator=() [1/2]

template<typename T>

CollectiveDeviceBuffer<T>& CollectiveDeviceBuffer< T >::operator= ( CollectiveDeviceBuffer< T > &&  other )

inline

Definition at line 88 of file CollectiveDeviceBuffer.h.

                                                                          {
    type = other.type;

    numElemsAlloc = other.numElemsAlloc;
    buffer = other.buffer;
    d_peerBuffers = other.d_peerBuffers;

    // This object owns the underlying data now, so we modify other
    other.numElemsAlloc = 0;
    other.buffer = nullptr;
    other.d_peerBuffers = nullptr;

    h_peerBuffers = std::move(other.h_peerBuffers);

    return *this;
  }

operator=() [2/2]

template<typename T>

CollectiveDeviceBuffer<T>& CollectiveDeviceBuffer< T >::operator= ( const CollectiveDeviceBuffer< T > &  other )

inline

Definition at line 106 of file CollectiveDeviceBuffer.h.

                                                                               {
    type = other.type;
    numElemsAlloc = other.numElemsAlloc;
    buffer = other.buffer;
    d_peerBuffers = other.d_peerBuffers;
    h_peerBuffers = other.h_peerBuffers;

    return *this;
  }

reallocate()

template<typename T >

void CollectiveDeviceBuffer< T >::reallocate	(	CollectiveBufferType	type_in,
		const size_t	newNumElems,
		const double	factor,
		SynchronousCollectiveScope	scope = SynchronousCollectiveScope::all
	)

Reallocates a symmetric device buffer on all devices if needed.

Definition at line 100 of file CollectiveDeviceBuffer.C.

References SynchronousCollectives::allReduce(), GlobalGPUMgr::getIsMasterPe(), GlobalGPUMgr::Object(), and SynchronousCollectives::Object().

                                    {
  SynchronousCollectives* syncColl = SynchronousCollectives::Object();
  GlobalGPUMgr* globalGPUMgr = GlobalGPUMgr::Object();
  const bool isMasterPe = globalGPUMgr->getIsMasterPe();

  size_t numElemRequested;
  if (isMasterPe) {
    numElemRequested = newNumElems;
  } else {
    numElemRequested = 0;
  }

  std::vector<size_t> numElemsVec = {numElemRequested};
  const size_t maxNumElemsRequested = syncColl->allReduce<size_t>(numElemsVec, CkReduction::max_ulong_long,
    scope)[0];

  reallocate_no_check(type_in, maxNumElemsRequested, factor);
}

reallocate_no_check()

template<typename T >

void CollectiveDeviceBuffer< T >::reallocate_no_check	(	CollectiveBufferType	type_in,
		const size_t	newNumElems,
		const double	factor
	)

Definition at line 121 of file CollectiveDeviceBuffer.C.

References Empty, GlobalGPUMgr::getIsMasterPe(), NAMD_die(), and GlobalGPUMgr::Object().

                       {

  GlobalGPUMgr* globalGPUMgr = GlobalGPUMgr::Object();
  const bool isMasterPe = globalGPUMgr->getIsMasterPe();
 
  if (isMasterPe) {
    if (type == CollectiveBufferType::Empty) {
      type = type_in;
    } else if (type != type_in) {
      NAMD_die("Reallocating buffer with different type");
    }

    const size_t maxNumElemsRequested = newNumElems;

    if (maxNumElemsRequested > numElemsAlloc) {
      const size_t newNumElemsAlloc = (size_t) ((double)maxNumElemsRequested * factor);
      deallocate();
      allocate_no_check(type_in, newNumElemsAlloc);
    }
  }
}

---

The documentation for this class was generated from the following files:

• CollectiveDeviceBuffer.h
• CollectiveDeviceBuffer.C