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 438 of file ComputePmeCUDAMgr.h.

Constructor & Destructor Documentation

ComputePmeCUDAMgr::ComputePmeCUDAMgr

(

)

Definition at line 204 of file ComputePmeCUDAMgr.C.

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

ComputePmeCUDAMgr::ComputePmeCUDAMgr

(

CkMigrateMessage *

)

Definition at line 215 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

(

)

Definition at line 227 of file ComputePmeCUDAMgr.C.

References cudaCheck, and stream.

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

Member Function Documentation

void ComputePmeCUDAMgr::activate_pencils

(

CkQdMsg *

msg

)

Definition at line 809 of file ComputePmeCUDAMgr.C.

References PmeStartMsg::data, and PmeStartMsg::dataSize.

                                                     {
  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();
    pmeStartYMsg->data = NULL;
    pmeStartYMsg->dataSize = 0;
    pmePencilY(ijPencilY[i].i, 0, ijPencilY[i].j).start(pmeStartYMsg);
  }

  for (int i=0;i < ijPencilZ.size();i++) {
    PmeStartMsg* pmeStartZMsg = new PmeStartMsg();
    pmeStartZMsg->data = NULL;
    pmeStartZMsg->dataSize = 0;
    pmePencilZ(ijPencilZ[i].i, ijPencilZ[i].j, 0).start(pmeStartZMsg);
  }

}

void ComputePmeCUDAMgr::createDevicesAndAtomFiler

(

)

Definition at line 679 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);

}

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

Definition at line 854 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;
}

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

Definition at line 895 of file ComputePmeCUDAMgr.C.

References getDevice().

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

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

Definition at line 903 of file ComputePmeCUDAMgr.C.

References getDevicePencilY().

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

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

Definition at line 911 of file ComputePmeCUDAMgr.C.

References getDevicePencilZ().

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

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

Definition at line 867 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;
}

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

Definition at line 882 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;
}

int ComputePmeCUDAMgr::getHomeNode

(

PatchID

patchID

)

Definition at line 845 of file ComputePmeCUDAMgr.C.

References getHomePencil(), and getNode().

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

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

Definition at line 238 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");
}

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

Definition at line 835 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;
}

PmeGrid ComputePmeCUDAMgr::getPmeGrid ( )

inline

Definition at line 448 of file ComputePmeCUDAMgr.h.

Referenced by ComputePmeCUDA::initialize().

{return pmeGrid;}

void ComputePmeCUDAMgr::initialize

(

CkQdMsg *

msg

)

Definition at line 598 of file ComputePmeCUDAMgr.C.

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

                                               {
        if (msg != NULL) delete msg;

  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));
    }
  }

}

void ComputePmeCUDAMgr::initialize_pencils

(

CkQdMsg *

msg

)

Definition at line 732 of file ComputePmeCUDAMgr.C.

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

                                                       {
  if (msg != NULL) delete msg;

  int numDevicesTmp = deviceCUDA->getNumDevice();

  // 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));
  }

  // 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));
  }

}

bool ComputePmeCUDAMgr::isPmeDevice

(

int

deviceID

)

Definition at line 585 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;
}

bool ComputePmeCUDAMgr::isPmeNode

(

int

node

)

Definition at line 573 of file ComputePmeCUDAMgr.C.

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

bool ComputePmeCUDAMgr::isPmePe

(

int

pe

)

Definition at line 563 of file ComputePmeCUDAMgr.C.

Referenced by CudaComputeNonbonded::assignPatches().

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

static ComputePmeCUDAMgr* ComputePmeCUDAMgr::Object ( )

inlinestatic

Definition at line 472 of file ComputePmeCUDAMgr.h.

Referenced by CudaComputeNonbonded::assignPatches().

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

void ComputePmeCUDAMgr::recvAtomFiler

(

CProxy_PmeAtomFiler

filer

)

Definition at line 698 of file ComputePmeCUDAMgr.C.

                                                               {
  pmeAtomFiler = filer;
}

void ComputePmeCUDAMgr::recvAtoms

(

PmeAtomMsg *

msg

)

Definition at line 933 of file ComputePmeCUDAMgr.C.

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

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

void ComputePmeCUDAMgr::recvDevices

(

RecvDeviceMsg *

msg

)

Definition at line 702 of file ComputePmeCUDAMgr.C.

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

                                                      {
  numDevicesMax = msg->numDevicesMax;
  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;
}

void ComputePmeCUDAMgr::recvPencils

(

CProxy_CudaPmePencilXYZ

xyz

)

Definition at line 713 of file ComputePmeCUDAMgr.C.

                                                               {
  pmePencilXYZ = xyz;
}

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

Definition at line 717 of file ComputePmeCUDAMgr.C.

References xy, and z.

                                                                                      {
  pmePencilXY = xy;
  pmePencilZ = z;
}

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

Definition at line 722 of file ComputePmeCUDAMgr.C.

References x, y, and z.

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

void ComputePmeCUDAMgr::setupPencils

(

)

Definition at line 400 of file ComputePmeCUDAMgr.C.

References PmeGrid::block1, PmeGrid::block2, PmeGrid::block3, deviceCUDA, PmeGrid::dim2, PmeGrid::dim3, DeviceCUDA::getNumDevice(), if(), PmeGrid::K1, PmeGrid::K2, PmeGrid::K3, NAMD_bug(), Node::Object(), PmeGrid::order, SimParameters::PMEGridSizeX, SimParameters::PMEGridSizeY, SimParameters::PMEGridSizeZ, SimParameters::PMEInterpOrder, SimParameters::PMEPencilsX, SimParameters::PMEPencilsY, SimParameters::PMEPencilsZ, Node::simParameters, PmeGrid::xBlocks, PmeGrid::yBlocks, and PmeGrid::zBlocks.

Referenced by initialize().

                                     {
  SimParameters *simParams = Node::Object()->simParameters;

  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);

  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;
  } else if (pmeGrid.yBlocks == 1) {
    // Blocks in all but y-dimension => 2D FFT
    pmePencilType = 2;
  } else {
    // Blocks in all dimensions => 1D FFT
    pmePencilType = 1;
  }

  //--------------------------------------------------------------------------
  // 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++;
    }
  }  
}

void ComputePmeCUDAMgr::skip

(

void

)

Definition at line 919 of file ComputePmeCUDAMgr.C.

                             {
  switch(pmePencilType) {
    case 1:
    pmePencilZ.skip();
    break;
    case 2:
    pmePencilZ.skip();
    break;
    case 3:
    pmePencilXYZ[0].skip();
    break;
  }
}

Member Data Documentation

ComputePmeCUDAMgr::ComputePmeCUDAMgr_SDAG_CODE

Definition at line 440 of file ComputePmeCUDAMgr.h.

---

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

• ComputePmeCUDAMgr.h
• ComputePmeCUDAMgr.C