ComputePmeCUDAMgr Class Reference

#include <ComputePmeCUDAMgr.h>

Inheritance diagram for ComputePmeCUDAMgr:

Public Member Functions
	ComputePmeCUDAMgr ()
	ComputePmeCUDAMgr (CkMigrateMessage *)
	~ComputePmeCUDAMgr ()
void	setupPencils ()
void	initialize (CkQdMsg *msg)
void	initialize_pencils (CkQdMsg *msg)
void	activate_pencils (CkQdMsg *msg)
PmeGrid	getPmeGrid ()
int	getNode (int i, int j)
int	getDevice (int i, int j)
int	getDevicePencilY (int i, int j)
int	getDevicePencilZ (int i, int j)
int	getDeviceIDPencilX (int i, int j)
int	getDeviceIDPencilY (int i, int j)
int	getDeviceIDPencilZ (int i, int j)
void	recvPencils (CProxy_CudaPmePencilXYZ xyz)
void	recvPencils (CProxy_CudaPmePencilXY xy, CProxy_CudaPmePencilZ z)
void	recvPencils (CProxy_CudaPmePencilX x, CProxy_CudaPmePencilY y, CProxy_CudaPmePencilZ z)
void	createDevicesAndAtomFiler ()
void	recvDevices (RecvDeviceMsg *msg)
void	recvAtomFiler (CProxy_PmeAtomFiler filer)
void	skip ()
void	recvAtoms (PmeAtomMsg *msg)
void	getHomePencil (PatchID patchID, int &homey, int &homez)
int	getHomeNode (PatchID patchID)
bool	isPmePe (int pe)
bool	isPmeNode (int node)
bool	isPmeDevice (int deviceID)

Static Public Member Functions
static ComputePmeCUDAMgr *	Object ()

Public Attributes
	ComputePmeCUDAMgr_SDAG_CODE

Detailed Description

Definition at line 579 of file ComputePmeCUDAMgr.h.

Constructor & Destructor Documentation

ComputePmeCUDAMgr() [1/2]

ComputePmeCUDAMgr::ComputePmeCUDAMgr

(

)

Definition at line 221 of file ComputePmeCUDAMgr.C.

                                     {
  __sdag_init();
  numDevices = 0;
  numTotalPatches = 0;
  numNodesContributed = 0;
  numDevicesMax = 0;
}

ComputePmeCUDAMgr() [2/2]

ComputePmeCUDAMgr::ComputePmeCUDAMgr

(

CkMigrateMessage *

)

Definition at line 232 of file ComputePmeCUDAMgr.C.

References NAMD_bug().

                                                       {
  __sdag_init();
  NAMD_bug("ComputePmeCUDAMgr cannot be migrated");
  numDevices = 0;
  numTotalPatches = 0;
  numNodesContributed = 0;
  numDevicesMax = 0;
}

~ComputePmeCUDAMgr()

ComputePmeCUDAMgr::~ComputePmeCUDAMgr

(

)

Definition at line 244 of file ComputePmeCUDAMgr.C.

References cudaCheck.

                                      {
  for (int i=0;i < extraDevices.size();i++) {
    cudaCheck(cudaSetDevice(extraDevices[i].deviceID));
    cudaCheck(cudaStreamDestroy(extraDevices[i].stream));
  }
}

Member Function Documentation

activate_pencils()

void ComputePmeCUDAMgr::activate_pencils

(

CkQdMsg *

msg

)

Definition at line 837 of file ComputePmeCUDAMgr.C.

References PmeStartMsg::dataGrid, PmeStartMsg::dataSizes, PmeStartMsg::enabledGrid, and NUM_GRID_MAX.

                                                     {
  if (msg != NULL) delete msg;

  for (int device=0;device < numDevices;device++) {
    deviceProxy[device].ckLocalBranch()->activate_pencils();
  }

  for (int i=0;i < ijPencilY.size();i++) {
    PmeStartMsg* pmeStartYMsg = new PmeStartMsg();
    for (unsigned iGrid = 0; iGrid < NUM_GRID_MAX; ++iGrid) {
      pmeStartYMsg->dataGrid[iGrid] = NULL;
      pmeStartYMsg->dataSizes[iGrid] = 0;
      pmeStartYMsg->enabledGrid[iGrid] = deviceProxy[0].ckLocalBranch()->isGridEnabled(iGrid);
    }
    pmePencilY(ijPencilY[i].i, 0, ijPencilY[i].j).start(pmeStartYMsg);
  }

  for (int i=0;i < ijPencilZ.size();i++) {
    PmeStartMsg* pmeStartZMsg = new PmeStartMsg();
    for (unsigned iGrid = 0; iGrid < NUM_GRID_MAX; ++iGrid) {
      pmeStartZMsg->dataGrid[iGrid] = NULL;
      pmeStartZMsg->dataSizes[iGrid] = 0;
      pmeStartZMsg->enabledGrid[iGrid] = deviceProxy[0].ckLocalBranch()->isGridEnabled(iGrid);
    }
    pmePencilZ(ijPencilZ[i].i, ijPencilZ[i].j, 0).start(pmeStartZMsg);
  }

}

createDevicesAndAtomFiler()

void ComputePmeCUDAMgr::createDevicesAndAtomFiler

(

)

Definition at line 705 of file ComputePmeCUDAMgr.C.

References RecvDeviceMsg::dev, NAMD_bug(), RecvDeviceMsg::numDevicesMax, and PRIORITY_SIZE.

                                                  {
  if (CkMyNode() != 0)
    NAMD_bug("ComputePmeCUDAMgr::createDevicesAndAtomFiler can only be called on root node");

  // Root node creates all device proxies
  // NOTE: Only root node has numDevicesMax
  RecvDeviceMsg* msg = new (numDevicesMax, PRIORITY_SIZE) RecvDeviceMsg();
  msg->numDevicesMax = numDevicesMax;
  for (int i=0;i < numDevicesMax;i++) {
    CProxy_ComputePmeCUDADevice dev = CProxy_ComputePmeCUDADevice::ckNew();
    memcpy(&msg->dev[i], &dev, sizeof(CProxy_ComputePmeCUDADevice));
  }
  thisProxy.recvDevices(msg);

  CProxy_PmeAtomFiler filer = CProxy_PmeAtomFiler::ckNew();
  thisProxy.recvAtomFiler(filer);

}

getDevice()

int ComputePmeCUDAMgr::getDevice	(	int	i,
		int	j
	)

Definition at line 888 of file ComputePmeCUDAMgr.C.

References NAMD_bug(), PmeGrid::yBlocks, and PmeGrid::zBlocks.

Referenced by getDeviceIDPencilX(), and recvAtoms().

                                             {
  if (i < 0 || i >= pmeGrid.yBlocks || j < 0 || j >= pmeGrid.zBlocks)
    NAMD_bug("ComputePmeCUDAMgr::getDevice, pencil index out of bounds");
  int ind = i + j*pmeGrid.yBlocks;
  int device = nodeDeviceList[ind].device;
  if (device == -1)
    NAMD_bug("ComputePmeCUDAMgr::getDevice, no device found");
  return device;
}

getDeviceIDPencilX()

int ComputePmeCUDAMgr::getDeviceIDPencilX	(	int	i,
		int	j
	)

Definition at line 929 of file ComputePmeCUDAMgr.C.

References getDevice().

                                                      {
  int device = getDevice(i, j);
  return deviceProxy[device].ckLocalBranch()->getDeviceID();
}

getDeviceIDPencilY()

int ComputePmeCUDAMgr::getDeviceIDPencilY	(	int	i,
		int	j
	)

Definition at line 937 of file ComputePmeCUDAMgr.C.

References getDevicePencilY().

                                                      {
  int device = getDevicePencilY(i, j);
  return deviceProxy[device].ckLocalBranch()->getDeviceID();
}

getDeviceIDPencilZ()

int ComputePmeCUDAMgr::getDeviceIDPencilZ	(	int	i,
		int	j
	)

Definition at line 945 of file ComputePmeCUDAMgr.C.

References getDevicePencilZ().

                                                      {
  int device = getDevicePencilZ(i, j);
  return deviceProxy[device].ckLocalBranch()->getDeviceID();
}

getDevicePencilY()

int ComputePmeCUDAMgr::getDevicePencilY	(	int	i,
		int	j
	)

Definition at line 901 of file ComputePmeCUDAMgr.C.

References NAMD_bug(), PmeGrid::xBlocks, and PmeGrid::zBlocks.

Referenced by getDeviceIDPencilY().

                                                    {
  if (i < 0 || i >= pmeGrid.xBlocks || j < 0 || j >= pmeGrid.zBlocks)
    NAMD_bug("ComputePmeCUDAMgr::getDevicePencilY, pencil index out of bounds");
  for (int device=0;device < ijPencilY.size();device++) {
    if (ijPencilY[device].i == i && ijPencilY[device].j == j) return device;
  }
  char str[256];
  sprintf(str, "ComputePmeCUDAMgr::getDevicePencilY, no device found at i %d j %d",i,j);
  NAMD_bug(str);
  return -1;
}

getDevicePencilZ()

int ComputePmeCUDAMgr::getDevicePencilZ	(	int	i,
		int	j
	)

Definition at line 916 of file ComputePmeCUDAMgr.C.

References NAMD_bug(), PmeGrid::xBlocks, and PmeGrid::yBlocks.

Referenced by getDeviceIDPencilZ().

                                                    {
  if (i < 0 || i >= pmeGrid.xBlocks || j < 0 || j >= pmeGrid.yBlocks)
    NAMD_bug("ComputePmeCUDAMgr::getDevicePencilZ, pencil index out of bounds");
  for (int device=0;device < ijPencilZ.size();device++) {
    if (ijPencilZ[device].i == i && ijPencilZ[device].j == j) return device;
  }
  NAMD_bug("ComputePmeCUDAMgr::getDevicePencilZ, no device found");
  return -1;
}

getHomeNode()

int ComputePmeCUDAMgr::getHomeNode

(

PatchID

patchID

)

Definition at line 879 of file ComputePmeCUDAMgr.C.

References getHomePencil(), and getNode().

                                                  {
  int homey, homez;
  getHomePencil(patchID, homey, homez);
  return getNode(homey, homez);
}

getHomePencil()

void ComputePmeCUDAMgr::getHomePencil	(	PatchID	patchID,
		int &	homey,
		int &	homez
	)

Definition at line 255 of file ComputePmeCUDAMgr.C.

References getPencilDim(), PmeGrid::K2, PmeGrid::K3, PatchMap::max_b(), PatchMap::max_c(), PatchMap::min_b(), PatchMap::min_c(), NAMD_bug(), PatchMap::Object(), PmeGrid::order, Perm_X_Y_Z, PmeGrid::yBlocks, and PmeGrid::zBlocks.

Referenced by getHomeNode().

                                                                             {
  PatchMap *patchMap = PatchMap::Object();

  BigReal miny = patchMap->min_b(patchID);
  BigReal maxy = patchMap->max_b(patchID);

  BigReal minz = patchMap->min_c(patchID);
  BigReal maxz = patchMap->max_c(patchID);

  // Determine home pencil = pencil with most overlap
  
  // Calculate patch grid coordinates
  int patch_y0 = floor((miny+0.5)*pmeGrid.K2);
  int patch_y1 = floor((maxy+0.5)*pmeGrid.K2)-1;
  int patch_z0 = floor((minz+0.5)*pmeGrid.K3);
  int patch_z1 = floor((maxz+0.5)*pmeGrid.K3)-1;

  if (patch_y0 < 0 || patch_y1 >= pmeGrid.K2 || patch_z0 < 0 || patch_z1 >= pmeGrid.K3) {
    NAMD_bug("ComputePmeCUDAMgr::getHomePencil, patch bounds are outside grid bounds");
  }

  int maxOverlap = 0;
  homey = -1;
  homez = -1;
  for (int iz=0;iz < pmeGrid.zBlocks;iz++) {
    for (int iy=0;iy < pmeGrid.yBlocks;iy++) {
      int pencil_x0, pencil_x1, pencil_y0, pencil_y1, pencil_z0, pencil_z1;
      getPencilDim(pmeGrid, Perm_X_Y_Z, iy, iz,
        pencil_x0, pencil_x1, pencil_y0, pencil_y1, pencil_z0, pencil_z1);

      if (pencil_y1 - pencil_y0 < pmeGrid.order || pencil_z1 - pencil_z0 < pmeGrid.order)
        NAMD_bug("ComputePmeCUDAMgr::getHomePencil, pencil size must be >= PMEInterpOrder");

      int y0 = std::max(patch_y0, pencil_y0);
      int y1 = std::min(patch_y1, pencil_y1);
      int z0 = std::max(patch_z0, pencil_z0);
      int z1 = std::min(patch_z1, pencil_z1);

      int overlap = (y1-y0 > 0 && z1-z0 > 0) ? (y1-y0)*(z1-z0) : -1;

      if (overlap > maxOverlap) {
        maxOverlap = overlap;
        homey = iy;
        homez = iz;
      }
    }
  }

  if (homey == -1 || homez == -1)
    NAMD_bug("ComputePmeCUDAMgr::getHomePencil, home pencil not found");
}

getNode()

int ComputePmeCUDAMgr::getNode	(	int	i,
		int	j
	)

Definition at line 869 of file ComputePmeCUDAMgr.C.

References NAMD_bug(), PmeGrid::yBlocks, and PmeGrid::zBlocks.

Referenced by getHomeNode().

                                           {
  if (i < 0 || i >= pmeGrid.yBlocks || j < 0 || j >= pmeGrid.zBlocks)
    NAMD_bug("ComputePmeCUDAMgr::getNode, pencil index out of bounds");
  int ind = i + j*pmeGrid.yBlocks;
  return nodeDeviceList[ind].node;
}

getPmeGrid()

PmeGrid ComputePmeCUDAMgr::getPmeGrid ( )

inline

Definition at line 589 of file ComputePmeCUDAMgr.h.

Referenced by ComputePmeCUDA::initialize().

{return pmeGrid;}

initialize()

void ComputePmeCUDAMgr::initialize

(

CkQdMsg *

msg

)

Definition at line 624 of file ComputePmeCUDAMgr.C.

References DebugM, endi(), iINFO(), iout, setupPencils(), PmeGrid::xBlocks, PmeGrid::yBlocks, and PmeGrid::zBlocks.

                                               {
        if (msg != NULL) delete msg;
  DebugM(4, "ComputePmeCUDAMgr::initialize(CkQdMsg)\n" << endi);
  setupPencils();

  if ( ! CkMyNode() ) {
    iout << iINFO << "PME using " << pmeGrid.xBlocks << " x " <<
      pmeGrid.yBlocks << " x " << pmeGrid.zBlocks <<
      " pencil grid for FFT and reciprocal sum.\n" << endi;
  }

  // Initialize list that contains the number of home patches for each device on this manager
  numHomePatchesList.resize(numDevices, 0);

  //--------------------------------------------------------------------------
  // Create devices and atom filer
  // numDevices = number of devices we'll be using, possibly different on each node
  // Use root node to compute the maximum number of devices in use over all nodes
  thisProxy[0].initializeDevicesAndAtomFiler(new NumDevicesMsg(numDevices));
  //--------------------------------------------------------------------------

  if (CkMyNode() == 0) {

    if (pmePencilType == 3) {
                // Single block => 3D FFT
      CProxy_PmePencilXYZMap xyzMap = CProxy_PmePencilXYZMap::ckNew(xPes[0]);
      CkArrayOptions xyzOpts(1);
      xyzOpts.setMap(xyzMap);
      xyzOpts.setAnytimeMigration(false);
      xyzOpts.setStaticInsertion(true);
      pmePencilXYZ = CProxy_CudaPmePencilXYZ::ckNew(xyzOpts);
      pmePencilXYZ[0].initialize(new CudaPmeXYZInitMsg(pmeGrid));
      thisProxy.recvPencils(pmePencilXYZ);
    } else if (pmePencilType == 2) {
      // Blocks in all but y-dimension => 2D FFT
      CProxy_PmePencilXYMap xyMap = CProxy_PmePencilXYMap::ckNew(xPes);
      CProxy_PmePencilXMap zMap = CProxy_PmePencilXMap::ckNew(0, 1, pmeGrid.xBlocks, zPes);
      CkArrayOptions xyOpts(1, 1, pmeGrid.zBlocks);
      CkArrayOptions zOpts(pmeGrid.xBlocks, 1, 1);
      xyOpts.setMap(xyMap);
      zOpts.setMap(zMap);
      xyOpts.setAnytimeMigration(false);
      zOpts.setAnytimeMigration(false);
      xyOpts.setStaticInsertion(true);
      zOpts.setStaticInsertion(true);
      pmePencilXY = CProxy_CudaPmePencilXY::ckNew(xyOpts);
      pmePencilZ = CProxy_CudaPmePencilZ::ckNew(zOpts);
      // Send pencil proxies to other nodes
      thisProxy.recvPencils(pmePencilXY, pmePencilZ);
      pmePencilXY.initialize(new CudaPmeXYInitMsg(pmeGrid, pmePencilXY, pmePencilZ, xyMap, zMap));
      pmePencilZ.initialize(new CudaPmeXYInitMsg(pmeGrid, pmePencilXY, pmePencilZ, xyMap, zMap));
    } else {
                // Blocks in all dimensions => 1D FFT
      CProxy_PmePencilXMap xMap = CProxy_PmePencilXMap::ckNew(1, 2, pmeGrid.yBlocks, xPes);
      CProxy_PmePencilXMap yMap = CProxy_PmePencilXMap::ckNew(0, 2, pmeGrid.xBlocks, yPes);
      CProxy_PmePencilXMap zMap = CProxy_PmePencilXMap::ckNew(0, 1, pmeGrid.xBlocks, zPes);
      CkArrayOptions xOpts(1, pmeGrid.yBlocks, pmeGrid.zBlocks);
      CkArrayOptions yOpts(pmeGrid.xBlocks, 1, pmeGrid.zBlocks);
                CkArrayOptions zOpts(pmeGrid.xBlocks, pmeGrid.yBlocks, 1);
      xOpts.setMap(xMap);
      yOpts.setMap(yMap);
      zOpts.setMap(zMap);
      xOpts.setAnytimeMigration(false);
      yOpts.setAnytimeMigration(false);
      zOpts.setAnytimeMigration(false);
      xOpts.setStaticInsertion(true);
      yOpts.setStaticInsertion(true);
      zOpts.setStaticInsertion(true);
      pmePencilX = CProxy_CudaPmePencilX::ckNew(xOpts);
      pmePencilY = CProxy_CudaPmePencilY::ckNew(yOpts);
      pmePencilZ = CProxy_CudaPmePencilZ::ckNew(zOpts);
      // Send pencil proxies to other nodes
      thisProxy.recvPencils(pmePencilX, pmePencilY, pmePencilZ);
      pmePencilX.initialize(new CudaPmeXInitMsg(pmeGrid, pmePencilX, pmePencilY, pmePencilZ, xMap, yMap, zMap));
      pmePencilY.initialize(new CudaPmeXInitMsg(pmeGrid, pmePencilX, pmePencilY, pmePencilZ, xMap, yMap, zMap));
      pmePencilZ.initialize(new CudaPmeXInitMsg(pmeGrid, pmePencilX, pmePencilY, pmePencilZ, xMap, yMap, zMap));
    }
  }

}

initialize_pencils()

void ComputePmeCUDAMgr::initialize_pencils

(

CkQdMsg *

msg

)

Definition at line 760 of file ComputePmeCUDAMgr.C.

References createStream(), cudaCheck, DebugM, deviceCUDA, endi(), DeviceCUDA::getDeviceIDbyRank(), and DeviceCUDA::getNumDevice().

                                                       {
  if (msg != NULL) delete msg;

  int numDevicesTmp = deviceCUDA->getNumDevice();
  DebugM(4, "ComputePmeCUDAMgr::initialize_pencils() numDevicesTmp " << numDevicesTmp <<"\n"<< endi);
  // Initialize device proxies for real-space interfacing
  for (int i=0;i < ijPencilX.size();i++) {
    // NOTE: i is here the device ID
    int deviceID = deviceCUDA->getDeviceIDbyRank(i % numDevicesTmp);
    deviceProxy[i].ckLocalBranch()->initialize(pmeGrid, ijPencilX[i].i, ijPencilX[i].j,
      deviceID, pmePencilType, thisProxy, pmeAtomFiler);
    if (pmePencilType == 1) {
      deviceProxy[i].ckLocalBranch()->setPencilProxy(pmePencilX);
    } else if (pmePencilType == 2) {
      deviceProxy[i].ckLocalBranch()->setPencilProxy(pmePencilXY);
    } else {
      deviceProxy[i].ckLocalBranch()->setPencilProxy(pmePencilXYZ);
    }
  }

  // Use above initialized device proxies for the PME pencils that interface with real-space
  for (int i=0;i < ijPencilX.size();i++) {
    if (pmePencilType == 1) {
      pmePencilX(0, ijPencilX[i].i, ijPencilX[i].j).initializeDevice(new InitDeviceMsg(deviceProxy[i]));
    } else if (pmePencilType == 2) {
      pmePencilXY(0, 0, ijPencilX[i].j).initializeDevice(new InitDeviceMsg(deviceProxy[i]));
    } else {
      pmePencilXYZ[0].initializeDevice(new InitDeviceMsg(deviceProxy[i]));
    }
  }

  // Create extra devices for Y and Z pencils if necessary
  int n = std::max(ijPencilY.size(), ijPencilZ.size());
  if (n > ijPencilX.size()) {
    int nextra = n - ijPencilX.size();
    extraDevices.resize(nextra);
    for (int i=0;i < nextra;i++) {
      extraDevices[i].deviceID = deviceCUDA->getDeviceIDbyRank((i + ijPencilX.size()) % numDevicesTmp);
      cudaCheck(cudaSetDevice(extraDevices[i].deviceID));
      createStream(extraDevices[i].stream);
    }
  }

  // Initialize Y pencils
  for (int i=0;i < ijPencilY.size();i++) {
    int deviceID;
    cudaStream_t stream;
    if (i < ijPencilX.size()) {
      deviceID = deviceProxy[i].ckLocalBranch()->getDeviceID();
      stream   = deviceProxy[i].ckLocalBranch()->getStream();
    } else {
      deviceID = extraDevices[i-ijPencilX.size()].deviceID;
      stream   = extraDevices[i-ijPencilX.size()].stream;
    }
    pmePencilY(ijPencilY[i].i, 0, ijPencilY[i].j).initializeDevice(new InitDeviceMsg2(deviceID, stream, thisProxy, deviceProxy[i]));
  }

  // Initialize Z pencils
  for (int i=0;i < ijPencilZ.size();i++) {
    int deviceID;
    cudaStream_t stream;
    if (i < ijPencilX.size()) {
      deviceID = deviceProxy[i].ckLocalBranch()->getDeviceID();
      stream   = deviceProxy[i].ckLocalBranch()->getStream();
    } else {
      deviceID = extraDevices[i-ijPencilX.size()].deviceID;
      stream   = extraDevices[i-ijPencilX.size()].stream;
    }
    pmePencilZ(ijPencilZ[i].i, ijPencilZ[i].j, 0).initializeDevice(new InitDeviceMsg2(deviceID, stream, thisProxy, deviceProxy[i]));
  }

}

isPmeDevice()

bool ComputePmeCUDAMgr::isPmeDevice

(

int

deviceID

)

Definition at line 611 of file ComputePmeCUDAMgr.C.

References deviceCUDA, DeviceCUDA::getDeviceIDbyRank(), and DeviceCUDA::getNumDevice().

                                                {
  for (int i=0;i < nodeDeviceList.size();i++) {
    if (deviceCUDA->getDeviceIDbyRank(nodeDeviceList[i].device % deviceCUDA->getNumDevice()) == deviceID) {
      return true;
    }
  }
  return false;
}

isPmeNode()

bool ComputePmeCUDAMgr::isPmeNode

(

int

node

)

Definition at line 599 of file ComputePmeCUDAMgr.C.

                                          {
  for (int i=0;i < nodeDeviceList.size();i++) {
    if (nodeDeviceList[i].node == node) {
      return true;
    }
  }
  return false;
}

isPmePe()

bool ComputePmeCUDAMgr::isPmePe

(

int

pe

)

Definition at line 589 of file ComputePmeCUDAMgr.C.

Referenced by CudaComputeNonbonded::assignPatches().

                                      {
  for (int i=0;i < xPes.size();i++) {
    if (pe == xPes[i]) return true;
  }
  return false;
}

Object()

static ComputePmeCUDAMgr* ComputePmeCUDAMgr::Object ( )

inlinestatic

Definition at line 613 of file ComputePmeCUDAMgr.h.

Referenced by CudaComputeNonbonded::assignPatches().

                                     {
    CProxy_ComputePmeCUDAMgr mgrProxy(CkpvAccess(BOCclass_group).computePmeCUDAMgr);
    return mgrProxy.ckLocalBranch();    
  }

recvAtomFiler()

void ComputePmeCUDAMgr::recvAtomFiler

(

CProxy_PmeAtomFiler

filer

)

Definition at line 724 of file ComputePmeCUDAMgr.C.

                                                               {
  pmeAtomFiler = filer;
}

recvAtoms()

void ComputePmeCUDAMgr::recvAtoms

(

PmeAtomMsg *

msg

)

Definition at line 968 of file ComputePmeCUDAMgr.C.

References DebugM, endi(), getDevice(), PmeAtomMsg::i, and PmeAtomMsg::j.

                                                 {
  DebugM(2, "ComputePmeCUDADevice::recvAtoms\n" << endi);
  int device = getDevice(msg->i, msg->j);
  deviceProxy[device].ckLocalBranch()->recvAtoms(msg);
}

recvDevices()

void ComputePmeCUDAMgr::recvDevices

(

RecvDeviceMsg *

msg

)

Definition at line 728 of file ComputePmeCUDAMgr.C.

References DebugM, RecvDeviceMsg::dev, endi(), NAMD_bug(), and RecvDeviceMsg::numDevicesMax.

                                                      {

  numDevicesMax = msg->numDevicesMax;
  DebugM(4, "ComputePmeCUDAMgr::recvDevices() numDevicesMax " << numDevicesMax <<"\n"<< endi);
  if (numDevices > numDevicesMax)
    NAMD_bug("ComputePmeCUDAMgr::recvDevices, numDevices > numDevicesMax");
  deviceProxy.resize(numDevices);
  for (int i=0;i < numDevices;i++) {
    deviceProxy[i] = msg->dev[i];
  }
  delete msg;
}

recvPencils() [1/3]

void ComputePmeCUDAMgr::recvPencils

(

CProxy_CudaPmePencilXYZ

xyz

)

Definition at line 741 of file ComputePmeCUDAMgr.C.

                                                               {
  pmePencilXYZ = xyz;
}

recvPencils() [2/3]

void ComputePmeCUDAMgr::recvPencils	(	CProxy_CudaPmePencilXY	xy,
		CProxy_CudaPmePencilZ	z
	)

Definition at line 745 of file ComputePmeCUDAMgr.C.

References xy.

                                                                                      {
  pmePencilXY = xy;
  pmePencilZ = z;
}

recvPencils() [3/3]

void ComputePmeCUDAMgr::recvPencils	(	CProxy_CudaPmePencilX	x,
		CProxy_CudaPmePencilY	y,
		CProxy_CudaPmePencilZ	z
	)

Definition at line 750 of file ComputePmeCUDAMgr.C.

                                                                                                             {
  pmePencilX = x;
  pmePencilY = y;
  pmePencilZ = z;
}

setupPencils()

void ComputePmeCUDAMgr::setupPencils

(

)

Definition at line 418 of file ComputePmeCUDAMgr.C.

References PmeGrid::block1, PmeGrid::block2, PmeGrid::block3, DebugM, deviceCUDA, PmeGrid::dim2, PmeGrid::dim3, endi(), DeviceCUDA::getNumDevice(), iINFO(), iout, PmeGrid::K1, PmeGrid::K2, PmeGrid::K3, NAMD_bug(), Node::Object(), PmeGrid::order, Node::simParameters, simParams, PmeGrid::xBlocks, PmeGrid::yBlocks, and PmeGrid::zBlocks.

Referenced by initialize().

                                     {
  SimParameters *simParams = Node::Object()->simParameters;
  DebugM(4, "ComputePmeCUDAMgr::setupPencils\n"<< endi);
  pmeGrid.K1 = simParams->PMEGridSizeX;
  pmeGrid.K2 = simParams->PMEGridSizeY;
  pmeGrid.K3 = simParams->PMEGridSizeZ;
  pmeGrid.order = simParams->PMEInterpOrder;
  pmeGrid.dim2 = pmeGrid.K2;
  pmeGrid.dim3 = 2 * (pmeGrid.K3/2 + 1);

  // Count the total number of devices assuming all nodes have the same number as this node
  // NOTE: This should be changed in the future to support heterogeneous nodes!!!
  int numDevicesTmp = deviceCUDA->getNumDevice();

  int numDeviceTot = CkNumNodes() * numDevicesTmp;
  // Use approximately 1/4th of the devices for PME
  // int numDeviceToUse = std::max(1, numDeviceTot/4);
  int numDeviceToUse = 1;// THERE CAN BE ONLY 1
  // there is a correctness problem in the multi-pencil case, so don't
  // go there until that is fixed and we have a better multi-device
  // parallelization.

  DebugM(4, "ComputePmeCUDAMgr::setupPencils numDeviceToUse "<<numDeviceToUse<< " numDeviceTot "<< numDeviceTot <<"\n"<< endi);
  if (numDeviceToUse < 4) {
    // 2D Slab
    pmeGrid.yBlocks = 1;
    pmeGrid.xBlocks = pmeGrid.zBlocks = numDeviceToUse;
  } else {
    // 1D Pencil
    pmeGrid.yBlocks = (int)sqrt((double)numDeviceToUse);
    pmeGrid.zBlocks = numDeviceToUse/pmeGrid.yBlocks;
    pmeGrid.xBlocks = pmeGrid.zBlocks;
  }

  if ( simParams->PMEPencilsX > 0 ) pmeGrid.xBlocks = simParams->PMEPencilsX;
  if ( simParams->PMEPencilsY > 0 ) pmeGrid.yBlocks = simParams->PMEPencilsY;
  if ( simParams->PMEPencilsZ > 0 ) pmeGrid.zBlocks = simParams->PMEPencilsZ;

  // Restrict number of pencils to the maximum number
  restrictToMaxPMEPencils();

  // Fix pencil numbers if they don't make sense w.r.t. number of devices
  if (pmeGrid.yBlocks == 1) {
    // 2D Slab
    if (pmeGrid.xBlocks > numDeviceTot) pmeGrid.xBlocks = numDeviceTot;
    if (pmeGrid.zBlocks > numDeviceTot) pmeGrid.zBlocks = numDeviceTot;
  } else {
    // 1D Pencil
    if (pmeGrid.yBlocks*pmeGrid.zBlocks > numDeviceTot ||
      pmeGrid.xBlocks*pmeGrid.zBlocks > numDeviceTot ||
      pmeGrid.xBlocks*pmeGrid.yBlocks > numDeviceTot) {
      pmeGrid.yBlocks = std::min(pmeGrid.yBlocks, (int)sqrt((double)numDeviceTot));
      pmeGrid.zBlocks = std::min(pmeGrid.zBlocks, numDeviceTot/pmeGrid.yBlocks);
    }
    pmeGrid.xBlocks = std::min(pmeGrid.yBlocks, pmeGrid.zBlocks);
  }

  // Here (block1, block2, block3) define the size of charge grid pencil in each direction
  pmeGrid.block1 = ( pmeGrid.K1 + pmeGrid.xBlocks - 1 ) / pmeGrid.xBlocks;
  pmeGrid.block2 = ( pmeGrid.K2 + pmeGrid.yBlocks - 1 ) / pmeGrid.yBlocks;
  pmeGrid.block3 = ( pmeGrid.K3 + pmeGrid.zBlocks - 1 ) / pmeGrid.zBlocks;

  // Determine type of FFT
  if (pmeGrid.xBlocks == 1 && pmeGrid.yBlocks == 1 && pmeGrid.zBlocks == 1) {
    // Single block => 3D FFT
    pmePencilType = 3;
    iout << iINFO << "Use 3D box decompostion in PME FFT.\n" << endi;
  } else if (pmeGrid.yBlocks == 1) {
    // Blocks in all but y-dimension => 2D FFT
    pmePencilType = 2;
    iout << iINFO << "Use 2D slab decompostion in PME FFT.\n" << endi;
  } else {
    // Blocks in all dimensions => 1D FFT
    pmePencilType = 1;
    iout << iINFO << "Use 1D pencil decompostion in PME FFT.\n" << endi;
  }

  //--------------------------------------------------------------------------
  // Map pencils into Pes
  xPes.resize(pmeGrid.yBlocks*pmeGrid.zBlocks);

  if (pmePencilType == 1 || pmePencilType == 2) {
    zPes.resize(pmeGrid.xBlocks*pmeGrid.yBlocks);
  }
  if (pmePencilType == 1) {
    yPes.resize(pmeGrid.xBlocks*pmeGrid.zBlocks);
  }  

    // i % numDeviceTot                              = device index
    // (i % numDeviceTot)/deviceCUDA->getNumDevice() = node index
    // (i % CkNodeSize(node))                        = pe displacement
  for (int i=0;i < xPes.size();i++) {
    int node = (i % numDeviceTot)/numDevicesTmp;
    xPes[i] = CkNodeFirst(node) + (i + 0) % CkNodeSize(node);
  }
  for (int i=0;i < yPes.size();i++) {
    int node = (i % numDeviceTot)/numDevicesTmp;
    yPes[i] = CkNodeFirst(node) + (i + 0) % CkNodeSize(node);
  }
  for (int i=0;i < zPes.size();i++) {
    int node = (i % numDeviceTot)/numDevicesTmp;
    zPes[i] = CkNodeFirst(node) + (i + 0) % CkNodeSize(node);
  }

  // char peStr[256];
  // char *p = peStr;
  // p += sprintf(p, "%2d | xPes", CkMyPe());
  // for (int i=0;i < xPes.size();i++)
  //   p += sprintf(p, " %d", xPes[i]);
  // p += sprintf(p, " yPes");
  // for (int i=0;i < yPes.size();i++)
  //   p += sprintf(p, " %d", yPes[i]);
  // p += sprintf(p, " zPes");
  // for (int i=0;i < zPes.size();i++)
  //   p += sprintf(p, " %d", zPes[i]);
  // fprintf(stderr, "%s | %d %d\n",peStr, CkNodeFirst(CkMyNode()), CkNodeSize(CkMyNode()));

  //--------------------------------------------------------------------------
  // Build global node list for x-pencils
  nodeDeviceList.resize(xPes.size());
  numDevices = 0;
  for (int k=0;k < xPes.size();k++) {
    nodeDeviceList[k].node = CkNodeOf(xPes[k]);
    nodeDeviceList[k].device = -1;
    if (nodeDeviceList[k].node == CkMyNode()) {
      nodeDeviceList[k].device = numDevices++;
    }
  }

  ijPencilX.clear();
  ijPencilY.clear();
  ijPencilZ.clear();

  // Construct list of pencil coordinates (i,j) for each device held by this node
  for (int k=0;k < xPes.size();k++) {
    if (CkMyNode() == CkNodeOf(xPes[k])) {
      IJ ij;
      ij.i = k % pmeGrid.yBlocks;
      ij.j = k / pmeGrid.yBlocks;
      ijPencilX.push_back(ij);
    }
  }
  if (ijPencilX.size() != numDevices)
    NAMD_bug("ComputePmeCUDAMgr::setupPencils, error setting up x-pencils and devices");

  int numDevicesY = 0;
  for (int k=0;k < yPes.size();k++) {
    if (CkMyNode() == CkNodeOf(yPes[k])) {
      IJ ij;
      ij.i = k % pmeGrid.xBlocks;
      ij.j = k / pmeGrid.xBlocks;
      ijPencilY.push_back(ij);
      numDevicesY++;
    }
  }

  int numDevicesZ = 0;
  for (int k=0;k < zPes.size();k++) {
    if (CkMyNode() == CkNodeOf(zPes[k])) {
      IJ ij;
      ij.i = k % pmeGrid.xBlocks;
      ij.j = k / pmeGrid.xBlocks;
      ijPencilZ.push_back(ij);
      numDevicesZ++;
    }
  }
}

skip()

void ComputePmeCUDAMgr::skip

(

void

)

Definition at line 953 of file ComputePmeCUDAMgr.C.

References DebugM, and endi().

                             {
  DebugM(2, "ComputePmeCUDADevice::skip\n" << endi);
  switch(pmePencilType) {
    case 1:
    pmePencilZ.skip();
    break;
    case 2:
    pmePencilZ.skip();
    break;
    case 3:
    pmePencilXYZ[0].skip();
    break;
  }
}

Member Data Documentation

ComputePmeCUDAMgr_SDAG_CODE

ComputePmeCUDAMgr::ComputePmeCUDAMgr_SDAG_CODE

Definition at line 581 of file ComputePmeCUDAMgr.h.

---

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

• ComputePmeCUDAMgr.h
• ComputePmeCUDAMgr.C