ComputePme.C File Reference

#include <sfftw.h>
#include <srfftw.h>
#include <vector>
#include <algorithm>
#include <deque>
#include "InfoStream.h"
#include "Node.h"
#include "PatchMap.h"
#include "PatchMap.inl"
#include "AtomMap.h"
#include "ComputePme.h"
#include "ComputePmeMgr.decl.h"
#include "PmeBase.inl"
#include "PmeRealSpace.h"
#include "PmeKSpace.h"
#include "ComputeNonbondedUtil.h"
#include "PatchMgr.h"
#include "Molecule.h"
#include "ReductionMgr.h"
#include "ComputeMgr.h"
#include "ComputeMgr.decl.h"
#include "Debug.h"
#include "SimParameters.h"
#include "WorkDistrib.h"
#include "varsizemsg.h"
#include "Random.h"
#include "ckhashtable.h"
#include "Priorities.h"
#include "CudaUtils.h"
#include "ComputeMoa.h"
#include "ComputeMoaMgr.decl.h"
#include "DeviceCUDA.h"
#include <cuda_runtime.h>
#include <cuda.h>
#include "ComputePmeCUDAKernel.h"
#include "ComputePmeMgr.def.h"

Go to the source code of this file.

Classes
class	PmeAckMsg
class	PmeGridMsg
class	PmeTransMsg
class	PmeSharedTransMsg
class	PmeUntransMsg
class	PmeSharedUntransMsg
class	PmeEvirMsg
class	PmePencilMap
struct	PmePencilInitMsgData
class	PmePencilInitMsg
struct	LocalPmeInfo
struct	NodePmeInfo
struct	sortop_bit_reversed
struct	ijpair
struct	ijpair_sortop_bit_reversed
class	ComputePmeMgr
struct	ComputePmeMgr::cuda_submit_charges_args
class	NodePmeMgr
class	PmePencil< T >
class	PmeZPencil
class	PmeYPencil
class	PmeXPencil

Macros
#define	fftwf_malloc   fftw_malloc
#define	fftwf_free   fftw_free
#define	MIN_DEBUG_LEVEL   3
#define	NUM_STREAMS   1
#define	CUDA_STREAM_CREATE(X)   cudaStreamCreate(X)
#define	CUDA_EVENT_ID_PME_CHARGES   80
#define	CUDA_EVENT_ID_PME_FORCES   81
#define	CUDA_EVENT_ID_PME_TICK   82
#define	CUDA_EVENT_ID_PME_COPY   83
#define	CUDA_EVENT_ID_PME_KERNEL   84
#define	count_limit   1000000
#define	CUDA_POLL(FN, ARG)   CcdCallFnAfter(FN,ARG,0.1)
#define	EVENT_STRIDE   10
#define	XCOPY(X)   masterPmeMgr->X = X;
#define	XCOPY(X)   X = masterPmeMgr->X;
#define	DEBUG_NODE_PAR_RECV   0

Functions
void	cuda_errcheck (const char *msg)
static int	findRecipEvirPe ()
void	generatePmePeList2 (int *gridPeMap, int numGridPes, int *transPeMap, int numTransPes)
int	compare_bit_reversed (int a, int b)
bool	less_than_bit_reversed (int a, int b)
ResizeArray< ComputePme * > &	getComputes (ComputePmeMgr *mgr)
int	isPmeProcessor (int p)
void	Pme_init ()
static void	PmeSlabSendTrans (int first, int last, void *result, int paraNum, void *param)
static void	PmeSlabSendUntrans (int first, int last, void *result, int paraNum, void *param)
static void	PmeSlabSendUngrid (int first, int last, void *result, int paraNum, void *param)
void	CcdCallBacksReset (void *ignored, double curWallTime)
void	cuda_check_pme_forces (void *arg, double walltime)
void	cuda_check_pme_charges (void *arg, double walltime)
static void	PmeXZPencilFFT (int first, int last, void *result, int paraNum, void *param)
static void	PmeZPencilSendTrans (int first, int last, void *result, int paraNum, void *param)
static void	PmeYPencilForwardFFT (int first, int last, void *result, int paraNum, void *param)
static void	PmeYPencilSendTrans (int first, int last, void *result, int paraNum, void *param)
static void	PmeXPencilSendUntrans (int first, int last, void *result, int paraNum, void *param)
static void	PmeYPencilBackwardFFT (int first, int last, void *result, int paraNum, void *param)
static void	PmeYPencilSendUntrans (int first, int last, void *result, int paraNum, void *param)
static void	PmeZPencilSendUngrid (int first, int last, void *result, int paraNum, void *param)

Variables
__thread DeviceCUDA *	deviceCUDA
char *	pencilPMEProcessors

Macro Definition Documentation

count_limit

#define count_limit   1000000

Definition at line 2504 of file ComputePme.C.

Referenced by cuda_check_pme_charges(), and cuda_check_pme_forces().

CUDA_EVENT_ID_PME_CHARGES

#define CUDA_EVENT_ID_PME_CHARGES   80

Referenced by ComputePmeMgr::ComputePmeMgr(), and cuda_check_pme_charges().

CUDA_EVENT_ID_PME_COPY

#define CUDA_EVENT_ID_PME_COPY   83

Referenced by ComputePmeMgr::chargeGridSubmitted(), ComputePmeMgr::ComputePmeMgr(), ComputePmeMgr::cuda_submit_charges(), and ComputePmeMgr::ungridCalc().

CUDA_EVENT_ID_PME_FORCES

#define CUDA_EVENT_ID_PME_FORCES   81

Referenced by ComputePmeMgr::ComputePmeMgr(), and cuda_check_pme_forces().

CUDA_EVENT_ID_PME_KERNEL

#define CUDA_EVENT_ID_PME_KERNEL   84

Referenced by ComputePmeMgr::ComputePmeMgr(), ComputePmeMgr::cuda_submit_charges(), and ComputePmeMgr::ungridCalc().

CUDA_EVENT_ID_PME_TICK

#define CUDA_EVENT_ID_PME_TICK   82

Referenced by ComputePmeMgr::ComputePmeMgr(), and ComputePmeMgr::ungridCalc().

CUDA_POLL

#define CUDA_POLL	(		FN,
			ARG
	)		CcdCallFnAfter(FN,ARG,0.1)

Definition at line 2505 of file ComputePme.C.

Referenced by cuda_check_pme_charges(), cuda_check_pme_forces(), ComputePmeMgr::pollChargeGridReady(), and ComputePmeMgr::pollForcesReady().

CUDA_STREAM_CREATE

#define CUDA_STREAM_CREATE

(

X

)

cudaStreamCreate(X)

Referenced by ComputePmeMgr::ComputePmeMgr().

DEBUG_NODE_PAR_RECV

#define DEBUG_NODE_PAR_RECV   0

Definition at line 5068 of file ComputePme.C.

EVENT_STRIDE

#define EVENT_STRIDE   10

Definition at line 2506 of file ComputePme.C.

Referenced by cuda_check_pme_forces(), and ComputePmeMgr::ungridCalc().

fftwf_free

#define fftwf_free   fftw_free

Definition at line 14 of file ComputePme.C.

Referenced by PmePencil< CBase_PmeZPencil >::~PmePencil().

fftwf_malloc

#define fftwf_malloc   fftw_malloc

Copyright (c) 1995, 1996, 1997, 1998, 1999, 2000 by The Board of Trustees of the University of Illinois. All rights reserved.

Definition at line 13 of file ComputePme.C.

Referenced by PmeZPencil::fft_init(), PmeYPencil::fft_init(), and PmeXPencil::fft_init().

MIN_DEBUG_LEVEL

#define MIN_DEBUG_LEVEL   3

Definition at line 47 of file ComputePme.C.

NUM_STREAMS

#define NUM_STREAMS   1

Definition at line 542 of file ComputePme.C.

Referenced by ComputePmeMgr::ComputePmeMgr().

XCOPY [1/2]

#define XCOPY

(

X

)

masterPmeMgr->X = X;

Referenced by ComputePmeMgr::initialize_computes().

XCOPY [2/2]

#define XCOPY

(

X

)

X = masterPmeMgr->X;

Function Documentation

CcdCallBacksReset()

void CcdCallBacksReset	(	void *	ignored,
		double	curWallTime
	)

Referenced by cuda_check_pme_charges(), cuda_check_pme_forces(), ComputePmeMgr::pollChargeGridReady(), and ComputePmeMgr::pollForcesReady().

compare_bit_reversed()

int compare_bit_reversed	(	int	a,
		int	b
	)

Definition at line 346 of file ComputePme.C.

                                       {
  int d = a ^ b;
  int c = 1;
  if ( d ) while ( ! (d & c) ) {
    c = c << 1;
  }
  return (a & c) - (b & c);
}

cuda_check_pme_charges()

void cuda_check_pme_charges	(	void *	arg,
		double	walltime
	)

Definition at line 3540 of file ComputePme.C.

References CcdCallBacksReset(), ComputePmeMgr::charges_time, ComputePmeMgr::check_charges_count, count_limit, CUDA_EVENT_ID_PME_CHARGES, CUDA_POLL, cudaDie(), ComputePmeMgr::end_charges, ComputePmeMgr::saved_sequence, and ComputePmeMgr::sendChargeGridReady().

Referenced by ComputePmeMgr::pollChargeGridReady().

                                                        {
  ComputePmeMgr *argp = (ComputePmeMgr *) arg;

  cudaError_t err = cudaEventQuery(argp->end_charges);
  if ( err == cudaSuccess ) {
    traceUserBracketEvent(CUDA_EVENT_ID_PME_CHARGES,argp->charges_time,walltime);
    argp->charges_time = walltime - argp->charges_time;
    argp->sendChargeGridReady();
    argp->check_charges_count = 0;
  } else if ( err != cudaErrorNotReady ) {
    char errmsg[256];
    sprintf(errmsg,"in cuda_check_pme_charges after polling %d times over %f s on seq %d",
            argp->check_charges_count, walltime - argp->charges_time,
            argp->saved_sequence);
    cudaDie(errmsg,err);
  } else if ( ++(argp->check_charges_count) >= count_limit ) {
    char errmsg[256];
    sprintf(errmsg,"cuda_check_pme_charges polled %d times over %f s on seq %d",
            argp->check_charges_count, walltime - argp->charges_time,
            argp->saved_sequence);
    cudaDie(errmsg,err);
  } else {
    CcdCallBacksReset(0,walltime);  // fix Charm++
    CUDA_POLL(cuda_check_pme_charges, arg);
  }
}

cuda_check_pme_forces()

void cuda_check_pme_forces	(	void *	arg,
		double	walltime
	)

Definition at line 2512 of file ComputePme.C.

References CcdCallBacksReset(), ComputePmeMgr::check_forces_count, count_limit, CUDA_EVENT_ID_PME_FORCES, CUDA_POLL, cudaDie(), ComputePmeMgr::end_forces, EVENT_STRIDE, ComputePmeMgr::forces_count, ComputePmeMgr::forces_done_count, ComputePmeMgr::forces_time, WorkDistrib::messageEnqueueWork(), ComputePmeMgr::pmeComputes, and ComputePmeMgr::saved_sequence.

Referenced by ComputePmeMgr::pollForcesReady().

                                                       {
  ComputePmeMgr *argp = (ComputePmeMgr *) arg;

 while ( 1 ) { // process multiple events per call
  cudaError_t err = cudaEventQuery(argp->end_forces[argp->forces_done_count/EVENT_STRIDE]);
  if ( err == cudaSuccess ) {
    argp->check_forces_count = 0;
    for ( int i=0; i<EVENT_STRIDE; ++i ) {
      WorkDistrib::messageEnqueueWork(argp->pmeComputes[argp->forces_done_count]);
      if ( ++(argp->forces_done_count) == argp->forces_count ) break;
    }
    if ( argp->forces_done_count == argp->forces_count ) { // last event
      traceUserBracketEvent(CUDA_EVENT_ID_PME_FORCES,argp->forces_time,walltime);
      argp->forces_time = walltime - argp->forces_time;
      //CkPrintf("cuda_check_pme_forces forces_time == %f\n", argp->forces_time);
      return;
    } else { // more events
      continue; // check next event
    }
  } else if ( err != cudaErrorNotReady ) {
    char errmsg[256];
    sprintf(errmsg,"in cuda_check_pme_forces for event %d after polling %d times over %f s on seq %d",
            argp->forces_done_count/EVENT_STRIDE,
            argp->check_forces_count, walltime - argp->forces_time,
            argp->saved_sequence);
    cudaDie(errmsg,err);
  } else if ( ++(argp->check_forces_count) >= count_limit ) {
    char errmsg[256];
    sprintf(errmsg,"cuda_check_pme_forces for event %d polled %d times over %f s on seq %d",
            argp->forces_done_count/EVENT_STRIDE,
            argp->check_forces_count, walltime - argp->forces_time,
            argp->saved_sequence);
    cudaDie(errmsg,err);
  } else {
    break; // call again
  }
 } // while ( 1 )
 CcdCallBacksReset(0,walltime);  // fix Charm++
 CUDA_POLL(cuda_check_pme_forces, arg);
}

cuda_errcheck()

void cuda_errcheck

(

const char *

msg

)

Definition at line 67 of file ComputePme.C.

References NAMD_die().

Referenced by ComputePmeMgr::ComputePmeMgr(), ComputePme::doWork(), ComputePmeMgr::initialize(), ComputePmeMgr::initialize_computes(), and ComputePmeMgr::ungridCalc().

                                    {
  cudaError_t err;
  if ((err = cudaGetLastError()) != cudaSuccess) {
    char host[128];
    gethostname(host, 128);  host[127] = 0;
    char devstr[128] = "";
    int devnum;
    if ( cudaGetDevice(&devnum) == cudaSuccess ) {
      sprintf(devstr, " device %d", devnum);
    }
    cudaDeviceProp deviceProp;
    if ( cudaGetDeviceProperties(&deviceProp, devnum) == cudaSuccess ) {
      sprintf(devstr, " device %d pci %x:%x:%x", devnum,
        deviceProp.pciDomainID, deviceProp.pciBusID, deviceProp.pciDeviceID);
    }
    char errmsg[1024];
    sprintf(errmsg,"CUDA error %s on Pe %d (%s%s): %s", msg, CkMyPe(), host, devstr, cudaGetErrorString(err));
    NAMD_die(errmsg);
  }
}

findRecipEvirPe()

static int findRecipEvirPe ( )

static

Definition at line 269 of file ComputePme.C.

References NAMD_bug(), PatchMap::numPatchesOnNode(), and PatchMap::Object().

Referenced by PmeXPencil::evir_init(), and ComputePmeMgr::initialize().

                             {
    PatchMap *patchMap = PatchMap::Object();
    {
      int mype = CkMyPe();
      if ( patchMap->numPatchesOnNode(mype) ) {
        return mype; 
      }
    }
    {
      int node = CmiMyNode();
      int firstpe = CmiNodeFirst(node);
      int nodeSize = CmiNodeSize(node);
      int myrank = CkMyRank();
      for ( int i=0; i<nodeSize; ++i ) {
        int pe = firstpe + (myrank+i)%nodeSize;
        if ( patchMap->numPatchesOnNode(pe) ) {
          return pe;
        }
      }
    }
    {
      int *pelist;
      int nodeSize;
      CmiGetPesOnPhysicalNode(CmiPhysicalNodeID(CkMyPe()), &pelist, &nodeSize);
      int myrank;
      for ( int i=0; i<nodeSize; ++i ) {
        if ( pelist[i] == CkMyPe() ) myrank = i;
      }
      for ( int i=0; i<nodeSize; ++i ) {
        int pe = pelist[(myrank+i)%nodeSize];
        if ( patchMap->numPatchesOnNode(pe) ) {
          return pe;
        }
      }
    }
    {
      int mype = CkMyPe();
      int npes = CkNumPes();
      for ( int i=0; i<npes; ++i ) {
        int pe = (mype+i)%npes;
        if ( patchMap->numPatchesOnNode(pe) ) {
          return pe;
        }
      }
    }
    NAMD_bug("findRecipEvirPe() failed!");
    return -999;  // should never happen
}

generatePmePeList2()

void generatePmePeList2	(	int *	gridPeMap,
		int	numGridPes,
		int *	transPeMap,
		int	numTransPes
	)

Definition at line 320 of file ComputePme.C.

References WorkDistrib::peDiffuseOrdering.

Referenced by ComputePmeMgr::initialize().

                                                                                         {
  int ncpus = CkNumPes();
  
  for ( int i=0; i<numGridPes; ++i ) {
    gridPeMap[i] = WorkDistrib::peDiffuseOrdering[ncpus - numGridPes + i];
  }
  std::sort(gridPeMap,gridPeMap+numGridPes);
  int firstTransPe = ncpus - numGridPes - numTransPes;
  if ( firstTransPe < 0 ) {
    firstTransPe = 0;
    // 0 should be first in list, skip if possible
    if ( ncpus > numTransPes ) firstTransPe = 1;
  }
  for ( int i=0; i<numTransPes; ++i ) {
    transPeMap[i] = WorkDistrib::peDiffuseOrdering[firstTransPe + i];
  }
  std::sort(transPeMap,transPeMap+numTransPes);
}

getComputes()

ResizeArray<ComputePme*>& getComputes

(

ComputePmeMgr *

mgr

)

Definition at line 615 of file ComputePme.C.

References ComputePmeMgr::pmeComputes.

Referenced by ComputeQM::saveResults().

                                                          {
    return mgr->pmeComputes ;
}

isPmeProcessor()

int isPmeProcessor

(

int

p

)

Definition at line 626 of file ComputePme.C.

References Node::Object(), pencilPMEProcessors, Node::simParameters, and simParams.

                         { 
  SimParameters *simParams = Node::Object()->simParameters;
  if (simParams->usePMECUDA) {
    return 0;
  } else {
    return pencilPMEProcessors[p];
  }
}

less_than_bit_reversed()

bool less_than_bit_reversed ( int  a, int  b  )

inline

Definition at line 355 of file ComputePme.C.

Referenced by sortop_bit_reversed::operator()(), and ijpair_sortop_bit_reversed::operator()().

                                                 {
  int d = a ^ b;
  int c = 1;
  if ( d ) while ( ! (d & c) ) {
    c = c << 1;
  }
  return d && (b & c);
}

Pme_init()

void Pme_init

(

)

Definition at line 882 of file ComputePme.C.

{
#if USE_TOPO_SFC
  if (CkMyRank() == 0) 
    tmgr_lock = CmiCreateLock();
#endif
}

PmeSlabSendTrans()

static void PmeSlabSendTrans ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 1962 of file ComputePme.C.

References ComputePmeMgr::sendTransSubset().

Referenced by ComputePmeMgr::sendTrans().

                                                                                                 {
  ComputePmeMgr *mgr = (ComputePmeMgr *)param;
  mgr->sendTransSubset(first, last);
}

PmeSlabSendUngrid()

static void PmeSlabSendUngrid ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 2399 of file ComputePme.C.

References ComputePmeMgr::sendUngridSubset().

Referenced by ComputePmeMgr::sendUngrid().

                                                                                                  {
  ComputePmeMgr *mgr = (ComputePmeMgr *)param;
  mgr->sendUngridSubset(first, last);
}

PmeSlabSendUntrans()

static void PmeSlabSendUntrans ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 2213 of file ComputePme.C.

References ComputePmeMgr::sendUntransSubset().

Referenced by ComputePmeMgr::sendUntrans().

                                                                                                   {
  ComputePmeMgr *mgr = (ComputePmeMgr *)param;
  mgr->sendUntransSubset(first, last);
}

PmeXPencilSendUntrans()

static void PmeXPencilSendUntrans ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 5830 of file ComputePme.C.

References PmeXPencil::send_subset_untrans().

Referenced by PmeXPencil::send_untrans().

                                                                                                     {
        PmeXPencil *xpencil = (PmeXPencil *)param;
        xpencil->send_subset_untrans(first, last);
}

PmeXZPencilFFT()

static void PmeXZPencilFFT ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 5266 of file ComputePme.C.

Referenced by PmeZPencil::backward_fft(), PmeXPencil::backward_fft(), PmeZPencil::forward_fft(), and PmeXPencil::forward_fft().

                                                                                              {
#ifdef NAMD_FFTW
#ifdef NAMD_FFTW_3    
    fftwf_plan *plans = (fftwf_plan *)param;
    for(int i=first; i<=last; i++) fftwf_execute(plans[i]);
#endif
#endif        
}

PmeYPencilBackwardFFT()

static void PmeYPencilBackwardFFT ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 6005 of file ComputePme.C.

References PmeYPencil::backward_subset_fft().

Referenced by PmeYPencil::backward_fft().

                                                                                                     {
        PmeYPencil *ypencil = (PmeYPencil *)param;
        ypencil->backward_subset_fft(first, last);
}

PmeYPencilForwardFFT()

static void PmeYPencilForwardFFT ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 5495 of file ComputePme.C.

References PmeYPencil::forward_subset_fft().

Referenced by PmeYPencil::forward_fft().

                                                                                                    {
        PmeYPencil *ypencil = (PmeYPencil *)param;
        ypencil->forward_subset_fft(first, last);
}

PmeYPencilSendTrans()

static void PmeYPencilSendTrans ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 5547 of file ComputePme.C.

References PmeYPencil::send_subset_trans().

Referenced by PmeYPencil::send_trans().

                                                                                                   {
        PmeYPencil *ypencil = (PmeYPencil *)param;
        ypencil->send_subset_trans(first, last);
}

PmeYPencilSendUntrans()

static void PmeYPencilSendUntrans ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 6064 of file ComputePme.C.

References PmeYPencil::send_subset_untrans().

Referenced by PmeYPencil::send_untrans().

                                                                                                     {
        PmeYPencil *ypencil = (PmeYPencil *)param;
        ypencil->send_subset_untrans(first, last);
}

PmeZPencilSendTrans()

static void PmeZPencilSendTrans ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 5331 of file ComputePme.C.

References PmeZPencil::send_subset_trans().

Referenced by PmeZPencil::send_trans().

                                                                                                   {
        PmeZPencil *zpencil = (PmeZPencil *)param;
        zpencil->send_subset_trans(first, last);        
}

PmeZPencilSendUngrid()

static void PmeZPencilSendUngrid ( int  first, int  last, void *  result, int  paraNum, void *  param  )

inlinestatic

Definition at line 6300 of file ComputePme.C.

References PmeZPencil::send_subset_ungrid().

Referenced by PmeZPencil::send_all_ungrid().

                                                                                                    {
        //to take advantage of the interface which allows 3 user params at most.
        //under such situtation, no new parameter list needs to be created!! -Chao Mei
        PmeZPencil *zpencil = (PmeZPencil *)param;
        zpencil->send_subset_ungrid(first, last);
}

Variable Documentation

deviceCUDA

__thread DeviceCUDA* deviceCUDA

Definition at line 23 of file DeviceCUDA.C.

Referenced by ComputePme::doWork(), ComputePmeMgr::initialize(), ComputePmeMgr::initialize_computes(), ComputePmeMgr::initialize_pencils(), and ComputePmeMgr::ungridCalc().

pencilPMEProcessors

char* pencilPMEProcessors

Definition at line 135 of file ComputePme.C.

Referenced by ComputePmeMgr::initialize(), and isPmeProcessor().