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 "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

Defines

#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 cudaDie (const char *msg, cudaError_t err=cudaSuccess)
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 DeviceCUDAdeviceCUDA
char * pencilPMEProcessors

Define Documentation

#define count_limit   1000000

Definition at line 2469 of file ComputePme.C.

#define CUDA_EVENT_ID_PME_CHARGES   80
#define CUDA_EVENT_ID_PME_COPY   83
#define CUDA_EVENT_ID_PME_FORCES   81
#define CUDA_EVENT_ID_PME_KERNEL   84
#define CUDA_EVENT_ID_PME_TICK   82
#define CUDA_POLL ( FN,
ARG   )     CcdCallFnAfter(FN,ARG,0.1)

Definition at line 2470 of file ComputePme.C.

#define CUDA_STREAM_CREATE (  )     cudaStreamCreate(X)
#define DEBUG_NODE_PAR_RECV   0

Definition at line 4953 of file ComputePme.C.

#define EVENT_STRIDE   10

Definition at line 2471 of file ComputePme.C.

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

#define fftwf_free   fftw_free

Definition at line 14 of file ComputePme.C.

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

#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 PmeXPencil::fft_init(), PmeYPencil::fft_init(), and PmeZPencil::fft_init().

#define MIN_DEBUG_LEVEL   3

Definition at line 47 of file ComputePme.C.

#define NUM_STREAMS   1

Definition at line 514 of file ComputePme.C.

Referenced by ComputePmeMgr::ComputePmeMgr().

#define XCOPY (  )     X = masterPmeMgr->X;
#define XCOPY (  )     masterPmeMgr->X = X;

Function Documentation

void CcdCallBacksReset ( void *  ignored,
double  curWallTime 
)
int compare_bit_reversed ( int  a,
int  b 
)

Definition at line 318 of file ComputePme.C.

Referenced by pe_sortop_bit_reversed::operator()().

00318                                        {
00319   int d = a ^ b;
00320   int c = 1;
00321   if ( d ) while ( ! (d & c) ) {
00322     c = c << 1;
00323   }
00324   return (a & c) - (b & c);
00325 }

void cuda_check_pme_charges ( void *  arg,
double  walltime 
)

Definition at line 3447 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().

03447                                                         {
03448   ComputePmeMgr *argp = (ComputePmeMgr *) arg;
03449 
03450   cudaError_t err = cudaEventQuery(argp->end_charges);
03451   if ( err == cudaSuccess ) {
03452     traceUserBracketEvent(CUDA_EVENT_ID_PME_CHARGES,argp->charges_time,walltime);
03453     argp->charges_time = walltime - argp->charges_time;
03454     argp->sendChargeGridReady();
03455     argp->check_charges_count = 0;
03456   } else if ( err != cudaErrorNotReady ) {
03457     char errmsg[256];
03458     sprintf(errmsg,"in cuda_check_pme_charges after polling %d times over %f s on seq %d",
03459             argp->check_charges_count, walltime - argp->charges_time,
03460             argp->saved_sequence);
03461     cudaDie(errmsg,err);
03462   } else if ( ++(argp->check_charges_count) >= count_limit ) {
03463     char errmsg[256];
03464     sprintf(errmsg,"cuda_check_pme_charges polled %d times over %f s on seq %d",
03465             argp->check_charges_count, walltime - argp->charges_time,
03466             argp->saved_sequence);
03467     cudaDie(errmsg,err);
03468   } else {
03469     CcdCallBacksReset(0,walltime);  // fix Charm++
03470     CUDA_POLL(cuda_check_pme_charges, arg);
03471   }
03472 }

void cuda_check_pme_forces ( void *  arg,
double  walltime 
)

Definition at line 2477 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().

02477                                                        {
02478   ComputePmeMgr *argp = (ComputePmeMgr *) arg;
02479 
02480  while ( 1 ) { // process multiple events per call
02481   cudaError_t err = cudaEventQuery(argp->end_forces[argp->forces_done_count/EVENT_STRIDE]);
02482   if ( err == cudaSuccess ) {
02483     argp->check_forces_count = 0;
02484     for ( int i=0; i<EVENT_STRIDE; ++i ) {
02485       WorkDistrib::messageEnqueueWork(argp->pmeComputes[argp->forces_done_count]);
02486       if ( ++(argp->forces_done_count) == argp->forces_count ) break;
02487     }
02488     if ( argp->forces_done_count == argp->forces_count ) { // last event
02489       traceUserBracketEvent(CUDA_EVENT_ID_PME_FORCES,argp->forces_time,walltime);
02490       argp->forces_time = walltime - argp->forces_time;
02491       //CkPrintf("cuda_check_pme_forces forces_time == %f\n", argp->forces_time);
02492       return;
02493     } else { // more events
02494       continue; // check next event
02495     }
02496   } else if ( err != cudaErrorNotReady ) {
02497     char errmsg[256];
02498     sprintf(errmsg,"in cuda_check_pme_forces for event %d after polling %d times over %f s on seq %d",
02499             argp->forces_done_count/EVENT_STRIDE,
02500             argp->check_forces_count, walltime - argp->forces_time,
02501             argp->saved_sequence);
02502     cudaDie(errmsg,err);
02503   } else if ( ++(argp->check_forces_count) >= count_limit ) {
02504     char errmsg[256];
02505     sprintf(errmsg,"cuda_check_pme_forces for event %d polled %d times over %f s on seq %d",
02506             argp->forces_done_count/EVENT_STRIDE,
02507             argp->check_forces_count, walltime - argp->forces_time,
02508             argp->saved_sequence);
02509     cudaDie(errmsg,err);
02510   } else {
02511     break; // call again
02512   }
02513  } // while ( 1 )
02514  CcdCallBacksReset(0,walltime);  // fix Charm++
02515  CUDA_POLL(cuda_check_pme_forces, arg);
02516 }

void cuda_errcheck ( const char *  msg  ) 

Definition at line 41 of file ComputeNonbondedCUDA.C.

References NAMD_die().

Referenced by ComputeNonbondedCUDA::ComputeNonbondedCUDA(), ComputePmeMgr::ComputePmeMgr(), cuda_bind_atom_params(), cuda_bind_atoms(), cuda_bind_exclusions(), cuda_bind_force_table(), cuda_bind_forces(), cuda_bind_GBIS_bornRad(), cuda_bind_GBIS_dEdaSum(), cuda_bind_GBIS_dHdrPrefix(), cuda_bind_GBIS_energy(), cuda_bind_GBIS_intRad(), cuda_bind_GBIS_psiSum(), cuda_bind_lj_table(), cuda_bind_patch_pairs(), cuda_bind_vdw_types(), cuda_bind_virials(), cuda_check_local_progress(), cuda_check_remote_progress(), cuda_GBIS_P1(), cuda_GBIS_P2(), cuda_GBIS_P3(), cuda_init(), cuda_nonbonded_forces(), ComputePme::doWork(), ComputeNonbondedCUDA::doWork(), ComputeNonbondedCUDA::finishReductions(), ComputePmeMgr::initialize(), ComputePmeMgr::initialize_computes(), and ComputePmeMgr::ungridCalc().

00041                                     {
00042   cudaError_t err;
00043   if ((err = cudaGetLastError()) != cudaSuccess) {
00044     char host[128];
00045     gethostname(host, 128);  host[127] = 0;
00046     char devstr[128] = "";
00047     int devnum;
00048     if ( cudaGetDevice(&devnum) == cudaSuccess ) {
00049       sprintf(devstr, " device %d", devnum);
00050     }
00051     cudaDeviceProp deviceProp;
00052     if ( cudaGetDeviceProperties(&deviceProp, devnum) == cudaSuccess ) {
00053       sprintf(devstr, " device %d pci %x:%x:%x", devnum,
00054         deviceProp.pciDomainID, deviceProp.pciBusID, deviceProp.pciDeviceID);
00055     }
00056     char errmsg[1024];
00057     sprintf(errmsg,"CUDA error %s on Pe %d (%s%s): %s", msg, CkMyPe(), host, devstr, cudaGetErrorString(err));
00058     NAMD_die(errmsg);
00059   }
00060 }

void cudaDie ( const char *  msg,
cudaError_t  err = cudaSuccess 
)

Definition at line 9 of file CudaUtils.C.

Referenced by cuda_check_local_progress(), cuda_check_pme_charges(), cuda_check_pme_forces(), cuda_check_progress(), cuda_check_remote_progress(), DeviceCUDA::initialize(), and read_CUDA_ARCH().

00009                                                {
00010   char host[128];
00011   gethostname(host, 128);  host[127] = 0;
00012   char devstr[128] = "";
00013   int devnum;
00014   if ( cudaGetDevice(&devnum) == cudaSuccess ) {
00015     sprintf(devstr, " device %d", devnum);
00016   }
00017   cudaDeviceProp deviceProp;
00018   if ( cudaGetDeviceProperties(&deviceProp, devnum) == cudaSuccess ) {
00019     sprintf(devstr, " device %d pci %x:%x:%x", devnum,
00020       deviceProp.pciDomainID, deviceProp.pciBusID, deviceProp.pciDeviceID);
00021   }
00022   char errmsg[1024];
00023   if (err == cudaSuccess) {
00024     sprintf(errmsg,"CUDA error %s on Pe %d (%s%s)", msg, CkMyPe(), host, devstr);
00025   } else {
00026     sprintf(errmsg,"CUDA error %s on Pe %d (%s%s): %s", msg, CkMyPe(), host, devstr, cudaGetErrorString(err));    
00027   }
00028   NAMD_die(errmsg);
00029 }

static int findRecipEvirPe (  )  [static]

Definition at line 241 of file ComputePme.C.

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

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

00241                              {
00242     PatchMap *patchMap = PatchMap::Object();
00243     {
00244       int mype = CkMyPe();
00245       if ( patchMap->numPatchesOnNode(mype) ) {
00246         return mype; 
00247       }
00248     }
00249     {
00250       int node = CmiMyNode();
00251       int firstpe = CmiNodeFirst(node);
00252       int nodeSize = CmiNodeSize(node);
00253       int myrank = CkMyRank();
00254       for ( int i=0; i<nodeSize; ++i ) {
00255         int pe = firstpe + (myrank+i)%nodeSize;
00256         if ( patchMap->numPatchesOnNode(pe) ) {
00257           return pe;
00258         }
00259       }
00260     }
00261     {
00262       int *pelist;
00263       int nodeSize;
00264       CmiGetPesOnPhysicalNode(CmiPhysicalNodeID(CkMyPe()), &pelist, &nodeSize);
00265       int myrank;
00266       for ( int i=0; i<nodeSize; ++i ) {
00267         if ( pelist[i] == CkMyPe() ) myrank = i;
00268       }
00269       for ( int i=0; i<nodeSize; ++i ) {
00270         int pe = pelist[(myrank+i)%nodeSize];
00271         if ( patchMap->numPatchesOnNode(pe) ) {
00272           return pe;
00273         }
00274       }
00275     }
00276     {
00277       int mype = CkMyPe();
00278       int npes = CkNumPes();
00279       for ( int i=0; i<npes; ++i ) {
00280         int pe = (mype+i)%npes;
00281         if ( patchMap->numPatchesOnNode(pe) ) {
00282           return pe;
00283         }
00284       }
00285     }
00286     NAMD_bug("findRecipEvirPe() failed!");
00287     return -999;  // should never happen
00288 }

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

Definition at line 292 of file ComputePme.C.

References WorkDistrib::peDiffuseOrdering, and sort.

Referenced by ComputePmeMgr::initialize().

00292                                                                                          {
00293   int ncpus = CkNumPes();
00294   
00295   for ( int i=0; i<numGridPes; ++i ) {
00296     gridPeMap[i] = WorkDistrib::peDiffuseOrdering[ncpus - numGridPes + i];
00297   }
00298   std::sort(gridPeMap,gridPeMap+numGridPes);
00299   int firstTransPe = ncpus - numGridPes - numTransPes;
00300   if ( firstTransPe < 0 ) {
00301     firstTransPe = 0;
00302     // 0 should be first in list, skip if possible
00303     if ( ncpus > numTransPes ) firstTransPe = 1;
00304   }
00305   for ( int i=0; i<numTransPes; ++i ) {
00306     transPeMap[i] = WorkDistrib::peDiffuseOrdering[firstTransPe + i];
00307   }
00308   std::sort(transPeMap,transPeMap+numTransPes);
00309 }

ResizeArray<ComputePme*>& getComputes ( ComputePmeMgr mgr  ) 

Definition at line 587 of file ComputePme.C.

References ComputePmeMgr::pmeComputes.

Referenced by ComputeQM::saveResults().

00587                                                           {
00588     return mgr->pmeComputes ;
00589 }

int isPmeProcessor ( int  p  ) 

Definition at line 598 of file ComputePme.C.

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

00598                          { 
00599   SimParameters *simParams = Node::Object()->simParameters;
00600   if (simParams->usePMECUDA) {
00601     return 0;
00602   } else {
00603     return pencilPMEProcessors[p];
00604   }
00605 }

bool less_than_bit_reversed ( int  a,
int  b 
) [inline]

Definition at line 327 of file ComputePme.C.

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

00327                                                  {
00328   int d = a ^ b;
00329   int c = 1;
00330   if ( d ) while ( ! (d & c) ) {
00331     c = c << 1;
00332   }
00333   return d && (b & c);
00334 }

void Pme_init (  ) 

Definition at line 854 of file ComputePme.C.

00855 {
00856 #if USE_TOPO_SFC
00857   if (CkMyRank() == 0) 
00858     tmgr_lock = CmiCreateLock();
00859 #endif
00860 }

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

Definition at line 1934 of file ComputePme.C.

References ComputePmeMgr::sendTransSubset().

Referenced by ComputePmeMgr::sendTrans().

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

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

Definition at line 2364 of file ComputePme.C.

References ComputePmeMgr::sendUngridSubset().

Referenced by ComputePmeMgr::sendUngrid().

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

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

Definition at line 2178 of file ComputePme.C.

References ComputePmeMgr::sendUntransSubset().

Referenced by ComputePmeMgr::sendUntrans().

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

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

Definition at line 5714 of file ComputePme.C.

References PmeXPencil::send_subset_untrans().

Referenced by PmeXPencil::send_untrans().

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

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

Definition at line 5151 of file ComputePme.C.

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

05151                                                                                               {
05152 #ifdef NAMD_FFTW
05153 #ifdef NAMD_FFTW_3    
05154     fftwf_plan *plans = (fftwf_plan *)param;
05155     for(int i=first; i<=last; i++) fftwf_execute(plans[i]);
05156 #endif
05157 #endif        
05158 }

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

Definition at line 5889 of file ComputePme.C.

References PmeYPencil::backward_subset_fft().

Referenced by PmeYPencil::backward_fft().

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

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

Definition at line 5380 of file ComputePme.C.

References PmeYPencil::forward_subset_fft().

Referenced by PmeYPencil::forward_fft().

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

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

Definition at line 5432 of file ComputePme.C.

References PmeYPencil::send_subset_trans().

Referenced by PmeYPencil::send_trans().

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

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

Definition at line 5948 of file ComputePme.C.

References PmeYPencil::send_subset_untrans().

Referenced by PmeYPencil::send_untrans().

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

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

Definition at line 5216 of file ComputePme.C.

References PmeZPencil::send_subset_trans().

Referenced by PmeZPencil::send_trans().

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

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

Definition at line 6184 of file ComputePme.C.

References PmeZPencil::send_subset_ungrid().

Referenced by PmeZPencil::send_all_ungrid().

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


Variable Documentation

Definition at line 18 of file DeviceCUDA.C.

Definition at line 107 of file ComputePme.C.

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


Generated on 16 Sep 2019 for NAMD by  doxygen 1.6.1