ComputeNonbondedCUDA Class Reference

#include <ComputeNonbondedCUDA.h>

Inheritance diagram for ComputeNonbondedCUDA:

List of all members.

Public Member Functions
	ComputeNonbondedCUDA (ComputeID c, ComputeMgr *mgr, ComputeNonbondedCUDA *m=0, int idx=-1)
	~ComputeNonbondedCUDA ()
void	atomUpdate ()
void	doWork ()
int	noWork ()
void	recvYieldDevice (int pe)
int	finishWork ()
void	messageFinishWork ()
void	build_exclusions ()
void	requirePatch (int pid)
void	assignPatches ()
void	registerPatches ()
Static Public Member Functions
void	build_lj_table ()
void	build_force_table ()
Public Attributes
LocalWorkMsg *	localWorkMsg2
int	workStarted
Lattice	lattice
int	doSlow
int	doEnergy
int	step
ResizeArray< int >	activePatches
ResizeArray< int >	localActivePatches
ResizeArray< int >	remoteActivePatches
ResizeArray< int >	hostedPatches
ResizeArray< int >	localHostedPatches
ResizeArray< int >	remoteHostedPatches
ResizeArray< patch_record >	patchRecords
ResizeArray< compute_record >	computeRecords
ResizeArray< compute_record >	localComputeRecords
ResizeArray< compute_record >	remoteComputeRecords
int	num_atom_records
int	num_local_atom_records
int	num_remote_atom_records
int	num_force_records
float4 *	forces
float4 *	slow_forces
GBReal *	psiSumH
GBReal *	dEdaSumH
PatchMap *	patchMap
AtomMap *	atomMap
SubmitReduction *	reduction
ComputeNonbondedCUDAKernel *	kernel
ComputeNonbondedCUDA *	master
int	masterPe
int	slaveIndex
ComputeNonbondedCUDA **	slaves
int *	slavePes
int	numSlaves

---

Constructor & Destructor Documentation

ComputeNonbondedCUDA::ComputeNonbondedCUDA (  ComputeID  c, ComputeMgr *  mgr, ComputeNonbondedCUDA *  m = 0, int  idx = -1 )

Definition at line 725 of file ComputeNonbondedCUDA.C. References atomMap, atomsChanged, build_exclusions(), computeMgr, computesChanged, cuda_init(), cudaCompute, end_local_download, end_remote_download, force_lists_ptr, localWorkMsg2, master, masterPe, NAMD_bug(), NAMD_die(), Node::Object(), AtomMap::Object(), PatchMap::Object(), patch_pairs_ptr, patchMap, SimParameters::pressureProfileOn, reduction, registerPatches(), Node::simParameters, slavePes, slaves, start_calc, and workStarted. : Compute(c), slaveIndex(idx) { #ifdef PRINT_GBIS CkPrintf("C.N.CUDA[%d]::constructor cid=%d\n", CkMyPe(), c); #endif // CkPrintf("create ComputeNonbondedCUDA\n"); master = m ? m : this; cudaCompute = this; computeMgr = mgr; patchMap = PatchMap::Object(); atomMap = AtomMap::Object(); reduction = 0; SimParameters *params = Node::Object()->simParameters; if (params->pressureProfileOn) { NAMD_die("pressure profile not supported in CUDA"); } atomsChanged = 1; computesChanged = 1; workStarted = 0; basePriority = PROXY_DATA_PRIORITY; localWorkMsg2 = new (PRIORITY_SIZE) LocalWorkMsg; if ( master != this ) { // I am slave masterPe = master->masterPe; master->slaves[slaveIndex] = this; if ( master->slavePes[slaveIndex] != CkMyPe() ) { NAMD_bug("ComputeNonbondedCUDA slavePes[slaveIndex] != CkMyPe"); } registerPatches(); return; } masterPe = CkMyPe(); cuda_init(); build_exclusions(); // cudaEventCreate(&start_upload); cudaEventCreateWithFlags(&start_calc,cudaEventDisableTiming); cudaEventCreateWithFlags(&end_remote_download,cudaEventDisableTiming); cudaEventCreateWithFlags(&end_local_download,cudaEventDisableTiming); patch_pairs_ptr = new ResizeArray<patch_pair>; force_lists_ptr = new ResizeArray<force_list>; }

ComputeNonbondedCUDA::~ComputeNonbondedCUDA (   )

Definition at line 773 of file ComputeNonbondedCUDA.C. { ; }

---

Member Function Documentation

void ComputeNonbondedCUDA::assignPatches (   )

Definition at line 851 of file ComputeNonbondedCUDA.C. References activePatches, computeMgr, ComputeNonbondedCUDA::patch_record::hostPe, j, PatchMap::node(), numSlaves, ReductionMgr::Object(), PatchMap::ObjectOnPe(), PatchMap::patch(), patchMap, patchRecords, pesSharingDevice, reduction, REDUCTIONS_BASIC, registerPatches(), ComputeMgr::sendCreateNonbondedCUDASlave(), ResizeArray< Elem >::size(), slavePes, slaves, and ReductionMgr::willSubmit(). Referenced by ComputeMgr::createComputes(). { int *pesOnNodeSharingDevice = new int[CkMyNodeSize()]; int numPesOnNodeSharingDevice = 0; int masterIndex = -1; for ( int i=0; i<numPesSharingDevice; ++i ) { int pe = pesSharingDevice[i]; if ( pe == CkMyPe() ) masterIndex = numPesOnNodeSharingDevice; if ( CkNodeOf(pe) == CkMyNode() ) { pesOnNodeSharingDevice[numPesOnNodeSharingDevice++] = pe; } } int npatches = activePatches.size(); if ( npatches ) { reduction = ReductionMgr::Object()->willSubmit(REDUCTIONS_BASIC); } int *count = new int[npatches]; memset(count, 0, sizeof(int)*npatches); int *pcount = new int[numPesOnNodeSharingDevice]; memset(pcount, 0, sizeof(int)*numPesOnNodeSharingDevice); int *rankpcount = new int[CkMyNodeSize()]; memset(rankpcount, 0, sizeof(int)*CkMyNodeSize()); char *table = new char[npatches*numPesOnNodeSharingDevice]; memset(table, 0, npatches*numPesOnNodeSharingDevice); int unassignedpatches = npatches; if ( 0 ) { // assign all to device pe for ( int i=0; i<npatches; ++i ) { int pid = activePatches[i]; patch_record &pr = patchRecords[pid]; pr.hostPe = CkMyPe(); } unassignedpatches = 0; pcount[masterIndex] = npatches; } else // assign if home pe and build table of natural proxies for ( int i=0; i<npatches; ++i ) { int pid = activePatches[i]; patch_record &pr = patchRecords[pid]; int homePe = patchMap->node(pid); for ( int j=0; j<numPesOnNodeSharingDevice; ++j ) { int pe = pesOnNodeSharingDevice[j]; if ( pe == homePe ) { pr.hostPe = pe; --unassignedpatches; pcount[j] += 1; } if ( PatchMap::ObjectOnPe(pe)->patch(pid) ) { table[i*numPesOnNodeSharingDevice+j] = 1; } } if ( pr.hostPe == -1 && CkNodeOf(homePe) == CkMyNode() ) { pr.hostPe = homePe; --unassignedpatches; rankpcount[CkRankOf(homePe)] += 1; } } // assign if only one pe has a required proxy int assignj = 0; for ( int i=0; i<npatches; ++i ) { int pid = activePatches[i]; patch_record &pr = patchRecords[pid]; if ( pr.hostPe != -1 ) continue; int c = 0; int lastj; for ( int j=0; j<numPesOnNodeSharingDevice; ++j ) { if ( table[i*numPesOnNodeSharingDevice+j] ) { ++c; lastj=j; } } count[i] = c; if ( c == 1 ) { pr.hostPe = pesOnNodeSharingDevice[lastj]; --unassignedpatches; pcount[lastj] += 1; } } while ( unassignedpatches ) { int i; for ( i=0; i<npatches; ++i ) { if ( ! table[i*numPesOnNodeSharingDevice+assignj] ) continue; int pid = activePatches[i]; patch_record &pr = patchRecords[pid]; if ( pr.hostPe != -1 ) continue; pr.hostPe = pesOnNodeSharingDevice[assignj]; --unassignedpatches; pcount[assignj] += 1; if ( ++assignj == numPesOnNodeSharingDevice ) assignj = 0; break; } if ( i<npatches ) continue; // start search again for ( i=0; i<npatches; ++i ) { int pid = activePatches[i]; patch_record &pr = patchRecords[pid]; if ( pr.hostPe != -1 ) continue; if ( count[i] ) continue; pr.hostPe = pesOnNodeSharingDevice[assignj]; --unassignedpatches; pcount[assignj] += 1; if ( ++assignj == numPesOnNodeSharingDevice ) assignj = 0; break; } if ( i<npatches ) continue; // start search again if ( ++assignj == numPesOnNodeSharingDevice ) assignj = 0; } for ( int i=0; i<npatches; ++i ) { int pid = activePatches[i]; patch_record &pr = patchRecords[pid]; // CkPrintf("Pe %d patch %d hostPe %d\n", CkMyPe(), pid, pr.hostPe); } slavePes = new int[CkMyNodeSize()]; slaves = new ComputeNonbondedCUDA*[CkMyNodeSize()]; numSlaves = 0; for ( int j=0; j<numPesOnNodeSharingDevice; ++j ) { int pe = pesOnNodeSharingDevice[j]; int rank = pe - CkNodeFirst(CkMyNode()); // CkPrintf("host %d sharing %d pe %d rank %d pcount %d rankpcount %d\n", // CkMyPe(),j,pe,rank,pcount[j],rankpcount[rank]); if ( pe == CkMyPe() ) continue; if ( ! pcount[j] && ! rankpcount[rank] ) continue; rankpcount[rank] = 0; // skip in rank loop below slavePes[numSlaves] = pe; computeMgr->sendCreateNonbondedCUDASlave(pe,numSlaves); ++numSlaves; } for ( int j=0; j<CkMyNodeSize(); ++j ) { int pe = CkNodeFirst(CkMyNode()) + j; // CkPrintf("host %d rank %d pe %d rankpcount %d\n", // CkMyPe(),j,pe,rankpcount[j]); if ( ! rankpcount[j] ) continue; if ( pe == CkMyPe() ) continue; slavePes[numSlaves] = pe; computeMgr->sendCreateNonbondedCUDASlave(pe,numSlaves); ++numSlaves; } registerPatches(); delete [] pesOnNodeSharingDevice; delete [] count; delete [] pcount; delete [] rankpcount; delete [] table; }

void ComputeNonbondedCUDA::atomUpdate (   )  [virtual]

Reimplemented from Compute. Definition at line 1103 of file ComputeNonbondedCUDA.C. References atomsChanged. { atomsChanged = 1; }

void ComputeNonbondedCUDA::build_exclusions (   )

Definition at line 536 of file ComputeNonbondedCUDA.C. References ResizeArray< Elem >::add(), ResizeArray< Elem >::begin(), cuda_bind_exclusions(), ResizeArray< Elem >::end(), exclusionsByAtom, ExclusionSignature::fullExclCnt, ExclusionSignature::fullOffset, Molecule::get_full_exclusions_for_atom(), ObjectArena< Type >::getNewArray(), int32, j, MAX_EXCLUSIONS, Node::molecule, NAMD_bug(), NAMD_die(), Molecule::numAtoms, Node::Object(), ResizeArray< Elem >::resize(), SET_EXCL, ResizeArray< Elem >::size(), and SortableResizeArray< Elem >::sort(). Referenced by ComputeNonbondedCUDA(). { Molecule *mol = Node::Object()->molecule; #ifdef MEM_OPT_VERSION int natoms = mol->exclSigPoolSize; #else int natoms = mol->numAtoms; #endif exclusionsByAtom = new int2[natoms]; // create unique sorted lists ObjectArena<int32> listArena; ResizeArray<int32*> unique_lists; int32 **listsByAtom = new int32*[natoms]; SortableResizeArray<int32> curList; for ( int i=0; i<natoms; ++i ) { curList.resize(0); curList.add(0); // always excluded from self #ifdef MEM_OPT_VERSION const ExclusionSignature *sig = mol->exclSigPool + i; int n = sig->fullExclCnt; for ( int j=0; j<n; ++j ) { curList.add(sig->fullOffset[j]); } n += 1; #else const int32 *mol_list = mol->get_full_exclusions_for_atom(i); int n = mol_list[0] + 1; for ( int j=1; j<n; ++j ) { curList.add(mol_list[j] - i); } #endif curList.sort(); int j; for ( j=0; j<unique_lists.size(); ++j ) { if ( n != unique_lists[j][0] ) continue; // no match int k; for ( k=0; k<n; ++k ) { if ( unique_lists[j][k+3] != curList[k] ) break; } if ( k == n ) break; // found match } if ( j == unique_lists.size() ) { // no match int32 *list = listArena.getNewArray(n+3); list[0] = n; int maxdiff = 0; maxdiff = -1 * curList[0]; if ( curList[n-1] > maxdiff ) maxdiff = curList[n-1]; list[1] = maxdiff; for ( int k=0; k<n; ++k ) { list[k+3] = curList[k]; } unique_lists.add(list); } listsByAtom[i] = unique_lists[j]; } // sort lists by maxdiff std::stable_sort(unique_lists.begin(), unique_lists.end(), exlist_sortop()); long int totalbits = 0; int nlists = unique_lists.size(); for ( int j=0; j<nlists; ++j ) { int32 *list = unique_lists[j]; int maxdiff = list[1]; list[2] = totalbits + maxdiff; totalbits += 2*maxdiff + 1; } for ( int i=0; i<natoms; ++i ) { exclusionsByAtom[i].x = listsByAtom[i][1]; // maxdiff exclusionsByAtom[i].y = listsByAtom[i][2]; // start } delete [] listsByAtom; if ( totalbits & 31 ) totalbits += ( 32 - ( totalbits & 31 ) ); if ( ! CkMyPe() ) { long int bytesneeded = totalbits / 8; CkPrintf("Info: Found %d unique exclusion lists needing %ld bytes\n", unique_lists.size(), bytesneeded); long int bytesavail = MAX_EXCLUSIONS * sizeof(unsigned int); if ( bytesneeded > bytesavail ) { char errmsg[512]; sprintf(errmsg,"%ld bytes of CUDA memory needed for exclusions " "but only %ld bytes can be addressed with 32-bit int.", bytesneeded, bytesavail); NAMD_die(errmsg); } } #define SET_EXCL(EXCL,BASE,DIFF) \ (EXCL)[((BASE)+(DIFF))>>5] |= (1<<(((BASE)+(DIFF))&31)) unsigned int *exclusion_bits = new unsigned int[totalbits/32]; memset(exclusion_bits, 0, totalbits/8); long int base = 0; for ( int i=0; i<unique_lists.size(); ++i ) { base += unique_lists[i][1]; if ( base != unique_lists[i][2] ) { NAMD_bug("ComputeNonbondedCUDA::build_exclusions base != stored"); } int n = unique_lists[i][0]; for ( int j=0; j<n; ++j ) { SET_EXCL(exclusion_bits,base,unique_lists[i][j+3]); } base += unique_lists[i][1] + 1; } cuda_bind_exclusions(exclusion_bits, totalbits/32); delete [] exclusion_bits; }

void ComputeNonbondedCUDA::build_force_table (   )  [static]

Definition at line 396 of file ComputeNonbondedCUDA.C. References BigReal, cuda_bind_force_table(), diffa, FORCE_TABLE_SIZE, int32, and table_i. Referenced by build_cuda_force_table(). { // static float4 t[FORCE_TABLE_SIZE]; float4 et[FORCE_TABLE_SIZE]; // energy table const BigReal r2_delta = ComputeNonbondedUtil:: r2_delta; const int r2_delta_exp = ComputeNonbondedUtil:: r2_delta_exp; // const int r2_delta_expc = 64 * (r2_delta_exp - 127); const int r2_delta_expc = 64 * (r2_delta_exp - 1023); double r2list[FORCE_TABLE_SIZE]; // double to match cpu code for ( int i=1; i<FORCE_TABLE_SIZE; ++i ) { double r = ((double) FORCE_TABLE_SIZE) / ( (double) i + 0.5 ); r2list[i] = r*r + r2_delta; } union { double f; int32 i[2]; } byte_order_test; byte_order_test.f = 1.0; // should occupy high-order bits only int32 *r2iilist = (int32*)r2list + ( byte_order_test.i[0] ? 0 : 1 ); for ( int i=1; i<FORCE_TABLE_SIZE; ++i ) { double r = ((double) FORCE_TABLE_SIZE) / ( (double) i + 0.5 ); int table_i = (r2iilist[2*i] >> 14) + r2_delta_expc; // table_i >= 0 if ( r > cutoff ) { t[i].x = 0.; t[i].y = 0.; t[i].z = 0.; t[i].w = 0.; et[i].x = 0.; et[i].y = 0.; et[i].z = 0.; et[i].w = 0.; continue; } BigReal diffa = r2list[i] - r2_table[table_i]; // coulomb 1/r or fast force // t[i].x = 1. / (r2 * r); // -1/r * d/dr r^-1 { BigReal table_a = fast_table[4*table_i]; BigReal table_b = fast_table[4*table_i+1]; BigReal table_c = fast_table[4*table_i+2]; BigReal table_d = fast_table[4*table_i+3]; BigReal grad = ( 3. * table_d * diffa + 2. * table_c ) * diffa + table_b; t[i].x = 2. * grad; BigReal ener = table_a + diffa * ( ( table_d * diffa + table_c ) * diffa + table_b); et[i].x = ener; } // pme correction for slow force // t[i].w = 0.; { BigReal table_a = scor_table[4*table_i]; BigReal table_b = scor_table[4*table_i+1]; BigReal table_c = scor_table[4*table_i+2]; BigReal table_d = scor_table[4*table_i+3]; BigReal grad = ( 3. * table_d * diffa + 2. * table_c ) * diffa + table_b; t[i].w = 2. * grad; BigReal ener = table_a + diffa * ( ( table_d * diffa + table_c ) * diffa + table_b); et[i].w = ener; } // vdw 1/r^6 // t[i].y = 6. / (r8); // -1/r * d/dr r^-6 { BigReal table_a = vdwb_table[4*table_i]; BigReal table_b = vdwb_table[4*table_i+1]; BigReal table_c = vdwb_table[4*table_i+2]; BigReal table_d = vdwb_table[4*table_i+3]; BigReal grad = ( 3. * table_d * diffa + 2. * table_c ) * diffa + table_b; t[i].y = 2. * -1. * grad; BigReal ener = table_a + diffa * ( ( table_d * diffa + table_c ) * diffa + table_b); et[i].y = -1. * ener; } // vdw 1/r^12 // t[i].z = 12e / (r8 * r4 * r2); // -1/r * d/dr r^-12 { BigReal table_a = vdwa_table[4*table_i]; BigReal table_b = vdwa_table[4*table_i+1]; BigReal table_c = vdwa_table[4*table_i+2]; BigReal table_d = vdwa_table[4*table_i+3]; BigReal grad = ( 3. * table_d * diffa + 2. * table_c ) * diffa + table_b; t[i].z = 2. * grad; BigReal ener = table_a + diffa * ( ( table_d * diffa + table_c ) * diffa + table_b); et[i].z = ener; } // CkPrintf("%d %g %g %g %g %g %g\n", i, r, diffa, // t[i].x, t[i].y, t[i].z, t[i].w); /* double r2 = r * r; double r4 = r2 * r2; double r8 = r4 * r4; t[i].x = 1. / (r2 * r); // -1/r * d/dr r^-1 t[i].y = 6. / (r8); // -1/r * d/dr r^-6 t[i].z = 12. / (r8 * r4 * r2); // -1/r * d/dr r^-12 t[i].w = 0.; */ } t[0].x = 0.f; t[0].y = 0.f; t[0].z = 0.f; t[0].w = 0.f; et[0].x = et[1].x; et[0].y = et[1].y; et[0].z = et[1].z; et[0].w = et[1].w; cuda_bind_force_table(t,et); if ( ! CkMyPe() ) { CkPrintf("Info: Updated CUDA force table with %d elements.\n", FORCE_TABLE_SIZE); } }

void ComputeNonbondedCUDA::build_lj_table (   )  [static]

Definition at line 369 of file ComputeNonbondedCUDA.C. References LJTable::TableEntry::A, LJTable::TableEntry::B, cuda_bind_lj_table(), LJTable::get_table_dim(), j, NAMD_bug(), and LJTable::table_val(). Referenced by build_cuda_force_table(). { // static const LJTable* const ljTable = ComputeNonbondedUtil:: ljTable; const int dim = ljTable->get_table_dim(); // round dim up to odd multiple of 16 int tsize = (((dim+16+31)/32)*32)-16; if ( tsize < dim ) NAMD_bug("ComputeNonbondedCUDA::build_lj_table bad tsize"); float2 *t = new float2[tsize*tsize]; float2 *row = t; for ( int i=0; i<dim; ++i, row += tsize ) { for ( int j=0; j<dim; ++j ) { const LJTable::TableEntry *e = ljTable->table_val(i,j); row[j].x = e->A * scaling; row[j].y = e->B * scaling; } } cuda_bind_lj_table(t,tsize); delete [] t; if ( ! CkMyPe() ) { CkPrintf("Info: Updated CUDA LJ table with %d x %d elements.\n", dim, dim); } }

void ComputeNonbondedCUDA::doWork (   )  [virtual]

Reimplemented from Compute. Definition at line 1189 of file ComputeNonbondedCUDA.C. References Lattice::a(), activePatches, atom_params, CompAtomExt::atomFixed, atoms, atomsChanged, Lattice::b(), ResizeArray< Elem >::begin(), ComputeNonbondedCUDA::patch_record::bornRad, bornRadH, Lattice::c(), PatchMap::center(), CompAtom::charge, computeRecords, computesChanged, COULOMB, cuda_bind_patch_pairs(), cuda_check_local_progress(), cuda_errcheck(), cudaCompute, dEdaSumH, ComputeNonbondedCUDA::patch_record::dHdrPrefix, dHdrPrefixH, ResizeArray< Elem >::end(), energy_gbis, exclusionsByAtom, finishWork(), Patch::flags, forces, SimParameters::GBISOn, GBISP, GBReal, Patch::getCudaAtomList(), Patch::getNumAtoms(), hostedPatches, CompAtomExt::id, ComputeNonbondedCUDA::patch_record::intRad, intRad0H, intRadSH, SubmitReduction::item(), j, kernel_launch_state, kernel_time, Flags::lattice, lattice, localActivePatches, localComputeRecords, ComputeNonbondedCUDA::patch_record::localStart, master, Flags::maxAtomMovement, Node::molecule, num_atom_records, num_atom_records_allocated, num_force_records, num_local_atom_records, num_remote_atom_records, num_virials, num_virials_allocated, ComputeNonbondedCUDA::patch_record::numAtoms, ComputeNonbondedCUDA::patch_record::numFreeAtoms, Node::Object(), ComputeNonbondedCUDA::compute_record::offset, SimParameters::outputCudaTiming, ComputeNonbondedCUDA::patch_record::p, pairlistsValid, Flags::pairlistTolerance, pairlistTolerance, Node::parameters, ComputeNonbondedCUDA::patch_record::patchID, patchMap, patchRecords, ComputeNonbondedCUDA::compute_record::pid, plcutoff2, CompAtom::position, psiSumH, recvYieldDevice(), reduction, REDUCTION_PAIRLIST_WARNINGS, remoteActivePatches, remoteComputeRecords, ResizeArray< Elem >::resize(), Flags::savePairlists, savePairlists, Compute::sequence(), Node::simParameters, simParams, ResizeArray< Elem >::size(), slow_forces, slow_virials, sortAtomsForCUDA(), CompAtomExt::sortOrder, Lattice::unscale(), Flags::usePairlists, usePairlists, CompAtom::vdwType, virials, workStarted, ComputeNonbondedCUDA::patch_record::x, Vector::x, ComputeNonbondedCUDA::patch_record::xExt, Vector::y, and Vector::z. { GBISP("C.N.CUDA[%d]::doWork: seq %d, phase %d, workStarted %d\n", \ CkMyPe(), sequence(), gbisPhase, workStarted); PATCH_PAIRS_REF; FORCE_LISTS_REF; if ( workStarted ) { //if work already started, check if finished if ( finishWork() ) { // finished workStarted = 0; basePriority = PROXY_DATA_PRIORITY; // higher to aid overlap } else { // need to call again workStarted = 2; basePriority = PROXY_RESULTS_PRIORITY; // lower for local if ( master == this && kernel_launch_state > 2 ) { cuda_check_local_progress(this,0.); // launches polling } } return; } workStarted = 1; basePriority = COMPUTE_PROXY_PRIORITY; Molecule *mol = Node::Object()->molecule; Parameters *params = Node::Object()->parameters; SimParameters *simParams = Node::Object()->simParameters; //execute only during GBIS phase 1, or if not using GBIS if (!simParams->GBISOn || gbisPhase == 1) { //bind new patches to GPU if ( atomsChanged || computesChanged ) { int npatches = activePatches.size(); pairlistsValid = 0; pairlistTolerance = 0.; if ( computesChanged ) { computesChanged = 0; num_virials = npatches; if ( num_virials > num_virials_allocated ) { if ( num_virials_allocated ) { cudaFreeHost(virials); cudaFreeHost(energy_gbis); cuda_errcheck("in cudaFreeHost virials"); } num_virials_allocated = 1.1 * num_virials + 1; cudaHostAlloc((void**)&virials,2*16*sizeof(float)*num_virials_allocated,cudaHostAllocMapped); cudaHostAlloc((void**)&energy_gbis,sizeof(float)*num_virials_allocated,cudaHostAllocMapped); for (int i = 0; i < num_virials_allocated; i++) energy_gbis[i] = 0.f; cuda_errcheck("in cudaHostAlloc virials"); slow_virials = virials + 16*num_virials; } int *ap = activePatches.begin(); for ( int i=0; i<localActivePatches.size(); ++i ) { *(ap++) = localActivePatches[i]; } for ( int i=0; i<remoteActivePatches.size(); ++i ) { *(ap++) = remoteActivePatches[i]; } // sort computes by distance between patches cr_sortop so(lattice); std::stable_sort(localComputeRecords.begin(),localComputeRecords.end(),so); std::stable_sort(remoteComputeRecords.begin(),remoteComputeRecords.end(),so); int nlc = localComputeRecords.size(); int nrc = remoteComputeRecords.size(); computeRecords.resize(nlc+nrc); compute_record *cr = computeRecords.begin(); for ( int i=0; i<nlc; ++i ) { *(cr++) = localComputeRecords[i]; } for ( int i=0; i<nrc; ++i ) { *(cr++) = remoteComputeRecords[i]; } force_lists.resize(npatches); for ( int i=0; i<npatches; ++i ) { patchRecords[activePatches[i]].localIndex = i; force_lists[i].force_list_size = 0; } int ncomputes = computeRecords.size(); patch_pairs.resize(ncomputes); for ( int i=0; i<ncomputes; ++i ) { ComputeNonbondedCUDA::compute_record &cr = computeRecords[i]; int lp1 = patchRecords[cr.pid[0]].localIndex; int lp2 = patchRecords[cr.pid[1]].localIndex; force_lists[lp1].force_list_size++; patch_pair &pp = patch_pairs[i]; pp.offset.x = cr.offset.x; pp.offset.y = cr.offset.y; pp.offset.z = cr.offset.z; } for ( int i=0; i<ncomputes; ++i ) { ComputeNonbondedCUDA::compute_record &cr = computeRecords[i]; int lp1 = patchRecords[cr.pid[0]].localIndex; patch_pair &pp = patch_pairs[i]; pp.patch1_force_list_index = lp1; pp.patch1_force_list_size = force_lists[lp1].force_list_size; } if ( 1 || simParams->outputCudaTiming ) { CkPrintf("Pe %d has %d local and %d remote patches and %d local and %d remote computes.\n", CkMyPe(), localActivePatches.size(), remoteActivePatches.size(), localComputeRecords.size(), remoteComputeRecords.size()); } } // computesChanged int istart = 0; int flstart = 0; int vlstart = 0; int max_atoms_per_patch = 0; int i; for ( i=0; i<npatches; ++i ) { if ( i == localActivePatches.size() ) { num_local_atom_records = istart; } force_lists[i].force_list_start = flstart; force_lists[i].force_output_start = istart; force_lists[i].atom_start = istart; force_lists[i].virial_list_start = vlstart; force_lists[i].virial_output_start = 16*i; patch_record &pr = patchRecords[activePatches[i]]; pr.localStart = istart; int natoms = pr.p->getNumAtoms(); int nfreeatoms = natoms; if ( fixedAtomsOn ) { const CompAtomExt *aExt = pr.xExt; for ( int j=0; j<natoms; ++j ) { if ( aExt[j].atomFixed ) --nfreeatoms; } } if ( natoms > max_atoms_per_patch ) max_atoms_per_patch = natoms; pr.numAtoms = natoms; pr.numFreeAtoms = nfreeatoms; force_lists[i].patch_size = nfreeatoms; if ( natoms & 15 ) { natoms += 16 - (natoms & 15); } if ( nfreeatoms & 15 ) { nfreeatoms += 16 - (nfreeatoms & 15); } force_lists[i].patch_stride = nfreeatoms; flstart += nfreeatoms * force_lists[i].force_list_size; vlstart += 16 * force_lists[i].force_list_size; istart += natoms; // already rounded up force_lists[i].force_list_size = 0; // rebuild below } if ( i == localActivePatches.size() ) { num_local_atom_records = istart; } num_force_records = flstart; num_atom_records = istart; num_remote_atom_records = num_atom_records - num_local_atom_records; if ( num_atom_records > num_atom_records_allocated ) { if ( num_atom_records_allocated ) { // delete [] atom_order; cudaFreeHost(atom_params); cudaFreeHost(atoms); cudaFreeHost(forces); cudaFreeHost(slow_forces); if (simParams->GBISOn) { cudaFreeHost(intRad0H);//6 GBIS arrays cudaFreeHost(intRadSH); cudaFreeHost(psiSumH); cudaFreeHost(bornRadH); cudaFreeHost(dEdaSumH); cudaFreeHost(dHdrPrefixH); } cuda_errcheck("in cudaFreeHost"); } num_atom_records_allocated = 1.1 * num_atom_records + 1; // atom_order = new int[num_atom_records_allocated]; cudaMallocHost((void**)&atom_params,sizeof(atom_param)*num_atom_records_allocated); cudaMallocHost((void**)&atoms,sizeof(atom)*num_atom_records_allocated); cuda_errcheck("in cudaMallocHost atoms"); cudaHostAlloc((void**)&forces,sizeof(float4)*num_atom_records_allocated,cudaHostAllocMapped); cudaHostAlloc((void**)&slow_forces,sizeof(float4)*num_atom_records_allocated,cudaHostAllocMapped); //allocate GBIS memory if (simParams->GBISOn) { cudaMallocHost((void**)&intRad0H,sizeof(float)*num_atom_records_allocated); cudaMallocHost((void**)&intRadSH,sizeof(float)*num_atom_records_allocated); cudaHostAlloc((void**)&psiSumH, sizeof(GBReal)*num_atom_records_allocated,cudaHostAllocMapped); cudaMallocHost((void**)&bornRadH,sizeof(float)*num_atom_records_allocated); cudaHostAlloc((void**)&dEdaSumH,sizeof(GBReal)*num_atom_records_allocated,cudaHostAllocMapped); cudaMallocHost((void**)&dHdrPrefixH,sizeof(float)*num_atom_records_allocated); } cuda_errcheck("in cudaHostAlloc forces"); } int bfstart = 0; int ncomputes = computeRecords.size(); for ( int i=0; i<ncomputes; ++i ) { ComputeNonbondedCUDA::compute_record &cr = computeRecords[i]; int p1 = cr.pid[0]; int p2 = cr.pid[1]; int lp1 = patchRecords[p1].localIndex; int lp2 = patchRecords[p2].localIndex; patch_pair &pp = patch_pairs[i]; pp.patch1_atom_start = patchRecords[p1].localStart; pp.patch1_force_start = force_lists[lp1].force_list_start + force_lists[lp1].patch_stride * force_lists[lp1].force_list_size; pp.virial_start = force_lists[lp1].virial_list_start + 16 * force_lists[lp1].force_list_size; pp.patch1_size = patchRecords[p1].numAtoms; pp.patch1_force_size = patchRecords[p1].numFreeAtoms; pp.patch2_force_size = patchRecords[p2].numFreeAtoms; pp.block_flags_start = bfstart; bfstart += ((pp.patch1_force_size + 127) >> 7) << 5; // istart += pp.patch1_size; // if ( istart & 15 ) { istart += 16 - (istart & 15); } pp.patch2_atom_start = patchRecords[p2].localStart; // pp.patch2_force_start = force_lists[lp2].force_list_start + // force_lists[lp2].patch_stride * force_lists[lp2].force_list_size; pp.patch2_size = patchRecords[p2].numAtoms; // pp.patch2_force_size = patchRecords[p2].numFreeAtoms; // this must happen at end to get self computes right force_lists[lp1].force_list_size++; // if ( lp1 != lp2 || cr.t[0] != cr.t[1] ) { // force_lists[lp2].force_list_size++; // } // istart += pp.patch2_size; // if ( istart & 15 ) { istart += 16 - (istart & 15); } } //for ncomputes #if 0 CkPrintf("Pe %d cuda_bind_patch_pairs %d %d %d %d %d\n", CkMyPe(), patch_pairs.size(), force_lists.size(), num_atom_records, num_force_records, max_atoms_per_patch); #endif int totalmem = patch_pairs.size() * sizeof(patch_pair) + force_lists.size() * sizeof(force_list) + num_force_records * sizeof(float4) + num_atom_records * sizeof(atom) + num_atom_records * sizeof(atom_param) + num_atom_records * sizeof(float4); int totalcopy = num_atom_records * ( sizeof(atom) + sizeof(float4) ); /* CkPrintf("Pe %d allocating %d MB of GPU memory, will copy %d kB per step\n", CkMyPe(), totalmem >> 20, totalcopy >> 10); */ cuda_bind_patch_pairs(patch_pairs.begin(), patch_pairs.size(), force_lists.begin(), force_lists.size(), num_atom_records, num_force_records, bfstart, max_atoms_per_patch); } // atomsChanged || computesChanged double charge_scaling = sqrt(COULOMB * scaling * dielectric_1); Flags &flags = patchRecords[hostedPatches[0]].p->flags; float maxAtomMovement = 0.; float maxPatchTolerance = 0.; for ( int i=0; i<activePatches.size(); ++i ) { patch_record &pr = patchRecords[activePatches[i]]; float maxMove = pr.p->flags.maxAtomMovement; if ( maxMove > maxAtomMovement ) maxAtomMovement = maxMove; float maxTol = pr.p->flags.pairlistTolerance; if ( maxTol > maxPatchTolerance ) maxPatchTolerance = maxTol; int start = pr.localStart; int n = pr.numAtoms; const CompAtom *a = pr.x; const CompAtomExt *aExt = pr.xExt; // int *ao = atom_order + start; if ( atomsChanged ) { #if 0 int nfree = pr.numFreeAtoms; if ( nfree != n ) { int k = 0; for ( int j=0; j<n; ++j ) { // put free atoms first if ( ! aExt[j].atomFixed ) ao[k++] = j; } if ( fixedAtomsOn ) for ( int j=0; j<n; ++j ) { // put fixed atoms at end if ( aExt[j].atomFixed ) ao[k++] = j; } } else { for ( int j=0; j<n; ++j ) { ao[j] = j; } } sortAtomsForCUDA(ao,a,nfree,n); #endif atom_param *ap = atom_params + start; for ( int k=0; k<n; ++k ) { int j = aExt[k].sortOrder; ap[k].vdw_type = a[j].vdwType; ap[k].index = aExt[j].id; #ifdef MEM_OPT_VERSION ap[k].excl_index = exclusionsByAtom[aExt[j].exclId].y; ap[k].excl_maxdiff = exclusionsByAtom[aExt[j].exclId].x; #else ap[k].excl_index = exclusionsByAtom[aExt[j].id].y; ap[k].excl_maxdiff = exclusionsByAtom[aExt[j].id].x; #endif } } { #if 1 const CudaAtom *ac = pr.p->getCudaAtomList(); atom *ap = atoms + start; memcpy(ap, ac, sizeof(atom)*n); #else Vector center = pr.p->flags.lattice.unscale(cudaCompute->patchMap->center(pr.patchID)); atom *ap = atoms + start; for ( int k=0; k<n; ++k ) { int j = aExt[k].sortOrder; ap[k].position.x = a[j].position.x - center.x; ap[k].position.y = a[j].position.y - center.y; ap[k].position.z = a[j].position.z - center.z; ap[k].charge = charge_scaling * a[j].charge; } #endif } } //GBISP("finished active patches\n") //CkPrintf("maxMovement = %f maxTolerance = %f save = %d use = %d\n", // maxAtomMovement, maxPatchTolerance, // flags.savePairlists, flags.usePairlists); savePairlists = 0; usePairlists = 0; if ( flags.savePairlists ) { savePairlists = 1; usePairlists = 1; } else if ( flags.usePairlists ) { if ( ! pairlistsValid || ( 2. * maxAtomMovement > pairlistTolerance ) ) { reduction->item(REDUCTION_PAIRLIST_WARNINGS) += 1; } else { usePairlists = 1; } } if ( ! usePairlists ) { pairlistsValid = 0; } float plcutoff = cutoff; if ( savePairlists ) { pairlistsValid = 1; pairlistTolerance = 2. * maxPatchTolerance; plcutoff += pairlistTolerance; } plcutoff2 = plcutoff * plcutoff; //CkPrintf("plcutoff = %f listTolerance = %f save = %d use = %d\n", // plcutoff, pairlistTolerance, savePairlists, usePairlists); #if 0 // calculate warp divergence if ( 1 ) { Flags &flags = patchRecords[hostedPatches[0]].p->flags; Lattice &lattice = flags.lattice; float3 lata, latb, latc; lata.x = lattice.a().x; lata.y = lattice.a().y; lata.z = lattice.a().z; latb.x = lattice.b().x; latb.y = lattice.b().y; latb.z = lattice.b().z; latc.x = lattice.c().x; latc.y = lattice.c().y; latc.z = lattice.c().z; int ncomputes = computeRecords.size(); for ( int ic=0; ic<ncomputes; ++ic ) { patch_pair &pp = patch_pairs[ic]; atom *a1 = atoms + pp.patch1_atom_start; int n1 = pp.patch1_size; atom *a2 = atoms + pp.patch2_atom_start; int n2 = pp.patch2_size; float offx = pp.offset.x * lata.x + pp.offset.y * latb.x + pp.offset.z * latc.x; float offy = pp.offset.x * lata.y + pp.offset.y * latb.y + pp.offset.z * latc.y; float offz = pp.offset.x * lata.z + pp.offset.y * latb.z + pp.offset.z * latc.z; // CkPrintf("%f %f %f\n", offx, offy, offz); int atoms_tried = 0; int blocks_tried = 0; int atoms_used = 0; int blocks_used = 0; for ( int ii=0; ii<n1; ii+=32 ) { // warps for ( int jj=0; jj<n2; jj+=16 ) { // shared atom loads int block_used = 0; for ( int j=jj; j<jj+16 && j<n2; ++j ) { // shared atoms int atom_used = 0; for ( int i=ii; i<ii+32 && i<n1; ++i ) { // threads float dx = offx + a1[i].position.x - a2[j].position.x; float dy = offy + a1[i].position.y - a2[j].position.y; float dz = offz + a1[i].position.z - a2[j].position.z; float r2 = dx*dx + dy*dy + dz*dz; if ( r2 < cutoff2 ) atom_used = 1; } ++atoms_tried; if ( atom_used ) { block_used = 1; ++atoms_used; } } ++blocks_tried; if ( block_used ) { ++blocks_used; } } } CkPrintf("blocks = %d/%d (%f) atoms = %d/%d (%f)\n", blocks_used, blocks_tried, blocks_used/(float)blocks_tried, atoms_used, atoms_tried, atoms_used/(float)atoms_tried); } } #endif } // !GBISOn || gbisPhase == 1 //Do GBIS if (simParams->GBISOn) { //open GBIS boxes depending on phase for ( int i=0; i<activePatches.size(); ++i ) { patch_record &pr = master->patchRecords[activePatches[i]]; GBISP("doWork[%d] accessing arrays for P%d\n",CkMyPe(),gbisPhase); if (gbisPhase == 1) { //Copy GBIS intRadius to Host if (atomsChanged) { float *intRad0 = intRad0H + pr.localStart; float *intRadS = intRadSH + pr.localStart; for ( int k=0; k<pr.numAtoms; ++k ) { int j = pr.xExt[k].sortOrder; intRad0[k] = pr.intRad[2*j+0]; intRadS[k] = pr.intRad[2*j+1]; } } } else if (gbisPhase == 2) { float *bornRad = bornRadH + pr.localStart; for ( int k=0; k<pr.numAtoms; ++k ) { int j = pr.xExt[k].sortOrder; bornRad[k] = pr.bornRad[j]; } } else if (gbisPhase == 3) { float *dHdrPrefix = dHdrPrefixH + pr.localStart; for ( int k=0; k<pr.numAtoms; ++k ) { int j = pr.xExt[k].sortOrder; dHdrPrefix[k] = pr.dHdrPrefix[j]; } } // end phases } // end for patches } // if GBISOn kernel_time = CkWallTimer(); kernel_launch_state = 1; if ( gpu_is_mine ) recvYieldDevice(-1); }

int ComputeNonbondedCUDA::finishWork (   )

Definition at line 1888 of file ComputeNonbondedCUDA.C. References ADD_TENSOR_OBJECT, atomsChanged, BigReal, ComputeNonbondedCUDA::patch_record::bornRad, ComputeNonbondedCUDA::patch_record::bornRadBox, Box< Owner, Data >::close(), cuda_errcheck(), cuda_timer_count, cuda_timer_total, ComputeNonbondedCUDA::patch_record::dEdaSum, ComputeNonbondedCUDA::patch_record::dEdaSumBox, dEdaSumH, ComputeNonbondedCUDA::patch_record::dHdrPrefix, ComputeNonbondedCUDA::patch_record::dHdrPrefixBox, end_local_download, end_remote_download, energy_gbis, Results::f, ComputeNonbondedCUDA::patch_record::f, Force, ComputeNonbondedCUDA::patch_record::forceBox, forces, SimParameters::GBISOn, GBISP, GBReal, Molecule::get_full_exclusions_for_atom(), ComputeNonbondedCUDA::patch_record::intRad, ComputeNonbondedCUDA::patch_record::intRadBox, SubmitReduction::item(), j, localHostedPatches, ComputeNonbondedCUDA::patch_record::localStart, master, mergegrids, Node::molecule, ComputeNonbondedCUDA::patch_record::numAtoms, ComputeNonbondedCUDA::patch_record::numFreeAtoms, Node::Object(), Box< Owner, Data >::open(), SimParameters::outputCudaTiming, ComputeNonbondedCUDA::patch_record::patchID, patchRecords, CompAtom::position, ComputeNonbondedCUDA::patch_record::positionBox, ComputeNonbondedCUDA::patch_record::psiSum, ComputeNonbondedCUDA::patch_record::psiSumBox, psiSumH, ComputeNonbondedCUDA::patch_record::r, reduction, REDUCTION_ELECT_ENERGY, REDUCTION_ELECT_ENERGY_SLOW, REDUCTION_LJ_ENERGY, REDUCTION_VIRIAL_NBOND, REDUCTION_VIRIAL_SLOW, remoteHostedPatches, Node::simParameters, simParams, ResizeArray< Elem >::size(), slow_forces, slow_virials, CompAtomExt::sortOrder, start_calc, step, SubmitReduction::submit(), virials, workStarted, Vector::x, ComputeNonbondedCUDA::patch_record::x, ComputeNonbondedCUDA::patch_record::xExt, Tensor::xx, Tensor::xy, Tensor::xz, Vector::y, Tensor::yx, Tensor::yy, Tensor::yz, Vector::z, Tensor::zx, Tensor::zy, and Tensor::zz. Referenced by doWork(). { GBISP("C.N.CUDA[%d]::fnWork: workStarted %d, phase %d\n", \ CkMyPe(), workStarted, gbisPhase) Molecule *mol = Node::Object()->molecule; SimParameters *simParams = Node::Object()->simParameters; ResizeArray<int> &patches( workStarted == 1 ? remoteHostedPatches : localHostedPatches ); //call once at beginning of each phase/step if ( !simParams->GBISOn || gbisPhase == 1 ) { for ( int i=0; i<patches.size(); ++i ) { patch_record &pr = master->patchRecords[patches[i]]; GBISP("GBIS[%d] fnWork() P0[%d] force.open()\n",CkMyPe(), pr.patchID); pr.r = pr.forceBox->open(); } } // !GBISOn || gbisPhase==1 // long long int wcount = 0; // double virial = 0.; // double virial_slow = 0.; for ( int i=0; i<patches.size(); ++i ) { // CkPrintf("Pe %d patch %d of %d pid %d\n",CkMyPe(),i,patches.size(),patches[i]); patch_record &pr = master->patchRecords[patches[i]]; int start = pr.localStart; const CompAtomExt *aExt = pr.xExt; if ( !simParams->GBISOn || gbisPhase == 3 ) { int nfree = pr.numFreeAtoms; pr.f = pr.r->f[Results::nbond]; Force *f = pr.f; Force *f_slow = pr.r->f[Results::slow]; const CompAtom *a = pr.x; // int *ao = atom_order + start; float4 *af = master->forces + start; float4 *af_slow = master->slow_forces + start; // only free atoms return forces for ( int k=0; k<nfree; ++k ) { int j = aExt[k].sortOrder; f[j].x += af[k].x; f[j].y += af[k].y; f[j].z += af[k].z; // wcount += af[k].w; // virial += af[k].w; if ( doSlow ) { f_slow[j].x += af_slow[k].x; f_slow[j].y += af_slow[k].y; f_slow[j].z += af_slow[k].z; // virial_slow += af_slow[k].w; } } #if 1 // check exclusions reported as w if ( CkNumPes() == 1 ) { const CompAtomExt *aExt = pr.xExt; for ( int k=0; k<nfree; ++k ) { int j = aExt[k].sortOrder; #ifdef MEM_OPT_VERSION int excl_expected = mol->exclSigPool[aExt[j].exclId].fullExclCnt + 1; #else int excl_expected = mol->get_full_exclusions_for_atom(aExt[j].id)[0] + 1; #endif if ( af[k].w != excl_expected ) { CkPrintf("%d:%d(%d) atom %d found %d exclusions but expected %d\n", i, j, k, aExt[j].id, (int)af[k].w, excl_expected ); } } } #endif } // !GBISOn || gbisPhase == 3 #if 0 if ( i % 31 == 0 ) for ( int j=0; j<3; ++j ) { CkPrintf("Pe %d patch %d atom %d (%f %f %f) force %f\n", CkMyPe(), i, j, pr.x[j].position.x, pr.x[j].position.y, pr.x[j].position.z, af[j].w); } #endif //Close Boxes depending on Phase if (simParams->GBISOn) { if (gbisPhase == 1) { //Copy dEdaSum from Host to Patch Box GBReal *psiSumMaster = master->psiSumH + start; for ( int k=0; k<pr.numAtoms; ++k ) { int j = aExt[k].sortOrder; pr.psiSum[j] += psiSumMaster[k]; } GBISP("C.N.CUDA[%d]::fnWork: P1 psiSum.close()\n", CkMyPe()); pr.psiSumBox->close(&(pr.psiSum)); } else if (gbisPhase == 2) { //Copy dEdaSum from Host to Patch Box GBReal *dEdaSumMaster = master->dEdaSumH + start; for ( int k=0; k<pr.numAtoms; ++k ) { int j = aExt[k].sortOrder; pr.dEdaSum[j] += dEdaSumMaster[k]; } GBISP("C.N.CUDA[%d]::fnWork: P2 dEdaSum.close()\n", CkMyPe()); pr.dEdaSumBox->close(&(pr.dEdaSum)); } else if (gbisPhase == 3) { GBISP("C.N.CUDA[%d]::fnWork: P3 all.close()\n", CkMyPe()); pr.intRadBox->close(&(pr.intRad)); //box 6 pr.bornRadBox->close(&(pr.bornRad)); //box 7 pr.dHdrPrefixBox->close(&(pr.dHdrPrefix)); //box 9 pr.positionBox->close(&(pr.x)); //box 0 pr.forceBox->close(&(pr.r)); } //end phases } else { //not GBIS GBISP("C.N.CUDA[%d]::fnWork: pos/force.close()\n", CkMyPe()); pr.positionBox->close(&(pr.x)); pr.forceBox->close(&(pr.r)); } }// end for #if 0 virial *= (-1./6.); reduction->item(REDUCTION_VIRIAL_NBOND_XX) += virial; reduction->item(REDUCTION_VIRIAL_NBOND_YY) += virial; reduction->item(REDUCTION_VIRIAL_NBOND_ZZ) += virial; if ( doSlow ) { virial_slow *= (-1./6.); reduction->item(REDUCTION_VIRIAL_SLOW_XX) += virial_slow; reduction->item(REDUCTION_VIRIAL_SLOW_YY) += virial_slow; reduction->item(REDUCTION_VIRIAL_SLOW_ZZ) += virial_slow; } #endif if ( workStarted == 1 && ! mergegrids && ( localHostedPatches.size() || master == this ) ) { GBISP("not finished, call again\n"); return 0; // not finished, call again } if ( master != this ) { // finished GBISP("finished\n"); if (simParams->GBISOn) gbisPhase = 1 + (gbisPhase % 3);//1->2->3->1... atomsChanged = 0; return 1; } cuda_timer_total += kernel_time; if ( !simParams->GBISOn || gbisPhase == 3 ) { { Tensor virial_tensor; BigReal energyv = 0.; BigReal energye = 0.; BigReal energys = 0.; for ( int i = 0; i < num_virials; ++i ) { virial_tensor.xx += virials[16*i]; virial_tensor.xy += virials[16*i+1]; virial_tensor.xz += virials[16*i+2]; virial_tensor.yx += virials[16*i+3]; virial_tensor.yy += virials[16*i+4]; virial_tensor.yz += virials[16*i+5]; virial_tensor.zx += virials[16*i+6]; virial_tensor.zy += virials[16*i+7]; virial_tensor.zz += virials[16*i+8]; energyv += virials[16*i+9]; energye += virials[16*i+10]; energys += virials[16*i+11]; if (simParams->GBISOn) { energye += energy_gbis[i]; } } ADD_TENSOR_OBJECT(reduction,REDUCTION_VIRIAL_NBOND,virial_tensor); if ( doEnergy ) { reduction->item(REDUCTION_LJ_ENERGY) += energyv; reduction->item(REDUCTION_ELECT_ENERGY) += energye; reduction->item(REDUCTION_ELECT_ENERGY_SLOW) += energys; } } if ( doSlow ) { Tensor virial_slow_tensor; for ( int i = 0; i < num_virials; ++i ) { virial_slow_tensor.xx += slow_virials[16*i]; virial_slow_tensor.xy += slow_virials[16*i+1]; virial_slow_tensor.xz += slow_virials[16*i+2]; virial_slow_tensor.yx += slow_virials[16*i+3]; virial_slow_tensor.yy += slow_virials[16*i+4]; virial_slow_tensor.yz += slow_virials[16*i+5]; virial_slow_tensor.zx += slow_virials[16*i+6]; virial_slow_tensor.zy += slow_virials[16*i+7]; virial_slow_tensor.zz += slow_virials[16*i+8]; } ADD_TENSOR_OBJECT(reduction,REDUCTION_VIRIAL_SLOW,virial_slow_tensor); } atomsChanged = 0; reduction->submit(); cuda_timer_count++; if ( simParams->outputCudaTiming && step % simParams->outputCudaTiming == 0 ) { // int natoms = mol->numAtoms; // double wpa = wcount; wpa /= natoms; // CkPrintf("Pe %d CUDA kernel %f ms, total %f ms, wpa %f\n", CkMyPe(), // kernel_time * 1.e3, time * 1.e3, wpa); #if 0 float upload_ms, remote_calc_ms; float local_calc_ms, total_ms; cuda_errcheck("before event timers"); cudaEventElapsedTime(&upload_ms, start_upload, start_calc); cuda_errcheck("in event timer 1"); cudaEventElapsedTime(&remote_calc_ms, start_calc, end_remote_download); cuda_errcheck("in event timer 2"); cudaEventElapsedTime(&local_calc_ms, end_remote_download, end_local_download); cuda_errcheck("in event timer 3"); cudaEventElapsedTime(&total_ms, start_upload, end_local_download); cuda_errcheck("in event timer 4"); cuda_errcheck("in event timers"); CkPrintf("CUDA EVENT TIMING: %d %f %f %f %f\n", CkMyPe(), upload_ms, remote_calc_ms, local_calc_ms, total_ms); #endif if ( cuda_timer_count >= simParams->outputCudaTiming ) { cuda_timer_total /= cuda_timer_count; CkPrintf("CUDA TIMING: %d %f ms/step on node %d\n", step, cuda_timer_total * 1.e3, CkMyPe()); } cuda_timer_count = 0; cuda_timer_total = 0; } } // !GBISOn || gbisPhase==3 // Next GBIS Phase GBISP("C.N.CUDA[%d]::fnWork: incrementing phase\n", CkMyPe()) if (simParams->GBISOn) gbisPhase = 1 + (gbisPhase % 3);//1->2->3->1... GBISP("C.N.CUDA[%d] finished ready for next step\n",CkMyPe()); return 1; // finished and ready for next step }

void ComputeNonbondedCUDA::messageFinishWork (   )

Definition at line 1875 of file ComputeNonbondedCUDA.C. References computeMgr, cuda_errcheck(), WorkDistrib::messageEnqueueWork(), Compute::priority(), ComputeMgr::sendNonbondedCUDASlaveEnqueue(), Compute::sequence(), slavePes, slaves, and workStarted. { cuda_errcheck("at cuda stream completed"); for ( int i = 0; i < numSlaves; ++i ) { computeMgr->sendNonbondedCUDASlaveEnqueue(slaves[i],slavePes[i],sequence(),priority(),workStarted); } WorkDistrib::messageEnqueueWork(this); }

int ComputeNonbondedCUDA::noWork (   )  [virtual]

Reimplemented from Compute. Definition at line 1121 of file ComputeNonbondedCUDA.C. References ComputeNonbondedCUDA::patch_record::bornRad, ComputeNonbondedCUDA::patch_record::bornRadBox, computeMgr, ComputeNonbondedCUDA::patch_record::dEdaSum, ComputeNonbondedCUDA::patch_record::dEdaSumBox, ComputeNonbondedCUDA::patch_record::dHdrPrefix, ComputeNonbondedCUDA::patch_record::dHdrPrefixBox, Flags::doEnergy, doEnergy, Flags::doFullElectrostatics, Flags::doNonbonded, doSlow, ComputeNonbondedCUDA::patch_record::forceBox, SimParameters::GBISOn, GBISP, Patch::getCompAtomExtInfo(), hostedPatches, ComputeNonbondedCUDA::patch_record::intRad, ComputeNonbondedCUDA::patch_record::intRadBox, Flags::lattice, lattice, master, masterPe, Node::Object(), Box< Owner, Data >::open(), ComputeNonbondedCUDA::patch_record::p, ComputeNonbondedCUDA::patch_record::patchID, patchRecords, ComputeNonbondedCUDA::patch_record::positionBox, ComputeNonbondedCUDA::patch_record::psiSum, ComputeNonbondedCUDA::patch_record::psiSumBox, reduction, ComputeMgr::sendNonbondedCUDASlaveReady(), Compute::sequence(), Node::simParameters, simParams, ResizeArray< Elem >::size(), Box< Owner, Data >::skip(), Flags::step, step, SubmitReduction::submit(), workStarted, ComputeNonbondedCUDA::patch_record::x, and ComputeNonbondedCUDA::patch_record::xExt. { SimParameters *simParams = Node::Object()->simParameters; Flags &flags = master->patchRecords[hostedPatches[0]].p->flags; lattice = flags.lattice; doSlow = flags.doFullElectrostatics; doEnergy = flags.doEnergy; step = flags.step; if ( ! flags.doNonbonded ) { GBISP("GBIS[%d] noWork() don't do nonbonded\n",CkMyPe()); for ( int i=0; i<hostedPatches.size(); ++i ) { patch_record &pr = master->patchRecords[hostedPatches[i]]; pr.positionBox->skip(); pr.forceBox->skip(); if (simParams->GBISOn) { pr.intRadBox->skip(); pr.psiSumBox->skip(); pr.bornRadBox->skip(); pr.dEdaSumBox->skip(); pr.dHdrPrefixBox->skip(); } } if ( master != this ) { computeMgr->sendNonbondedCUDASlaveReady(masterPe, hostedPatches.size(),atomsChanged,sequence()); } if ( reduction ) reduction->submit(); return 1; } for ( int i=0; i<hostedPatches.size(); ++i ) { patch_record &pr = master->patchRecords[hostedPatches[i]]; if (!simParams->GBISOn || gbisPhase == 1) { GBISP("GBIS[%d] noWork() P0[%d] open()\n",CkMyPe(), pr.patchID); pr.x = pr.positionBox->open(); pr.xExt = pr.p->getCompAtomExtInfo(); } if (simParams->GBISOn) { if (gbisPhase == 1) { GBISP("GBIS[%d] noWork() P1[%d] open()\n",CkMyPe(),pr.patchID); pr.intRad = pr.intRadBox->open(); pr.psiSum = pr.psiSumBox->open(); } else if (gbisPhase == 2) { GBISP("GBIS[%d] noWork() P2[%d] open()\n",CkMyPe(),pr.patchID); pr.bornRad = pr.bornRadBox->open(); pr.dEdaSum = pr.dEdaSumBox->open(); } else if (gbisPhase == 3) { GBISP("GBIS[%d] noWork() P3[%d] open()\n",CkMyPe(),pr.patchID); pr.dHdrPrefix = pr.dHdrPrefixBox->open(); } GBISP("opened GBIS boxes"); } } if ( master == this ) return 0; //work to do, enqueue as usual // message masterPe computeMgr->sendNonbondedCUDASlaveReady(masterPe, hostedPatches.size(),atomsChanged,sequence()); workStarted = 1; basePriority = COMPUTE_PROXY_PRIORITY; return 1; }

void ComputeNonbondedCUDA::recvYieldDevice (  int  pe  )

Definition at line 1680 of file ComputeNonbondedCUDA.C. References Lattice::a(), SimParameters::alpha_cutoff, atom_params, atoms, Lattice::b(), bornRadH, Lattice::c(), SimParameters::coulomb_radius_offset, cuda_bind_atom_params(), cuda_bind_atoms(), 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_virials(), cuda_check_local_calc(), cuda_check_local_progress(), cuda_check_remote_calc(), cuda_check_remote_progress(), cuda_GBIS_P1(), cuda_GBIS_P2(), cuda_GBIS_P3(), cuda_nonbonded_forces(), CUDA_POLL, SimParameters::cutoff, dEdaSumH, dHdrPrefixH, SimParameters::dielectric, end_local_download, end_remote_download, energy_gbis, forces, SimParameters::fsMax, SimParameters::GBISOn, GBISP, gpu_is_mine, intRad0H, intRadSH, SimParameters::kappa, kernel_launch_state, lattice, localActivePatches, localComputeRecords, SimParameters::nonbondedScaling, Node::Object(), plcutoff2, psiSumH, remote_submit_time, remoteActivePatches, remoteComputeRecords, Compute::sequence(), shared_gpu, Node::simParameters, simParams, ResizeArray< Elem >::size(), slow_forces, SimParameters::solvent_dielectric, start_calc, stream, stream2, SimParameters::switchingActive, SimParameters::switchingDist, virials, workStarted, Vector::x, Vector::y, and Vector::z. Referenced by doWork(), and ComputeMgr::recvYieldDevice(). { GBISP("C.N.CUDA[%d]::recvYieldDevice: seq %d, workStarted %d, \ gbisPhase %d, kls %d, from pe %d\n", CkMyPe(), sequence(), \ workStarted, gbisPhase, kernel_launch_state, pe) float3 lata, latb, latc; lata.x = lattice.a().x; lata.y = lattice.a().y; lata.z = lattice.a().z; latb.x = lattice.b().x; latb.y = lattice.b().y; latb.z = lattice.b().z; latc.x = lattice.c().x; latc.y = lattice.c().y; latc.z = lattice.c().z; SimParameters *simParams = Node::Object()->simParameters; switch ( kernel_launch_state ) { // Remote case 1: GBISP("C.N.CUDA[%d]::recvYieldDeviceR: case 1\n", CkMyPe()) ++kernel_launch_state; gpu_is_mine = 0; remote_submit_time = CkWallTimer(); if (!simParams->GBISOn || gbisPhase == 1) { // cudaEventRecord(start_upload, stream); if ( atomsChanged ) { cuda_bind_atom_params(atom_params); } if ( simParams->GBISOn) { cuda_bind_GBIS_psiSum(psiSumH); if ( atomsChanged ) { cuda_bind_GBIS_intRad(intRad0H, intRadSH); } } cuda_bind_atoms(atoms); cuda_bind_forces(forces, slow_forces); cuda_bind_virials(virials); cuda_bind_GBIS_energy(energy_gbis); cudaEventRecord(start_calc, stream); //call CUDA Kernels cuda_nonbonded_forces(lata, latb, latc, cutoff2, plcutoff2, localComputeRecords.size(),remoteComputeRecords.size(), localActivePatches.size(),remoteActivePatches.size(), doSlow, doEnergy, usePairlists, savePairlists, stream); if (simParams->GBISOn) { cuda_GBIS_P1( localComputeRecords.size(),remoteComputeRecords.size(), localActivePatches.size(),remoteActivePatches.size(), simParams->alpha_cutoff-simParams->fsMax, simParams->coulomb_radius_offset, lata, latb, latc, stream); } //cuda_load_forces(forces, (doSlow ? slow_forces : 0 ), // num_local_atom_records,num_remote_atom_records); cudaEventRecord(end_remote_download, stream); CUDA_POLL(cuda_check_remote_progress,this); if ( shared_gpu && ! mergegrids ) { CUDA_POLL(cuda_check_remote_calc,this); break; } } // !GBIS or gbisPhase==1 if (simParams->GBISOn) { if (gbisPhase == 1) { //GBIS P1 Kernel launched in previous code block } else if (gbisPhase == 2) { GBISP("C.N.CUDA[%d]::recvYieldDeviceR: <<<P2>>>\n", CkMyPe()) // cudaEventRecord(start_upload, stream); cuda_bind_GBIS_bornRad(bornRadH); cuda_bind_GBIS_dEdaSum(dEdaSumH); cudaEventRecord(start_calc, stream); cuda_GBIS_P2( localComputeRecords.size(),remoteComputeRecords.size(), localActivePatches.size(),remoteActivePatches.size(), (simParams->alpha_cutoff-simParams->fsMax), simParams->cutoff, simParams->nonbondedScaling, simParams->kappa, (simParams->switchingActive ? simParams->switchingDist : -1.0), simParams->dielectric, simParams->solvent_dielectric, lata, latb, latc, doEnergy, doSlow, stream ); cudaEventRecord(end_remote_download, stream); CUDA_POLL(cuda_check_remote_progress,this); } else if (gbisPhase == 3) { GBISP("C.N.CUDA[%d]::recvYieldDeviceR: <<<P3>>>\n", CkMyPe()) // cudaEventRecord(start_upload, stream); cuda_bind_GBIS_dHdrPrefix(dHdrPrefixH); cudaEventRecord(start_calc, stream); if (doSlow) cuda_GBIS_P3( localComputeRecords.size(),remoteComputeRecords.size(), localActivePatches.size(),remoteActivePatches.size(), (simParams->alpha_cutoff-simParams->fsMax), simParams->coulomb_radius_offset, simParams->nonbondedScaling, lata, latb, latc, stream ); cudaEventRecord(end_remote_download, stream); CUDA_POLL(cuda_check_remote_progress,this); } } // Local case 2: GBISP("C.N.CUDA[%d]::recvYieldDeviceL: case 2\n", CkMyPe()) ++kernel_launch_state; gpu_is_mine = 0; #if 0 cudaStreamWaitEvent(stream2, start_calc, 0); #else #define stream2 stream #endif if (!simParams->GBISOn || gbisPhase == 1) { cuda_nonbonded_forces(lata, latb, latc, cutoff2, plcutoff2, 0,localComputeRecords.size(), 0,localActivePatches.size(), doSlow, doEnergy, usePairlists, savePairlists, stream2); if (simParams->GBISOn) { cuda_GBIS_P1( 0,localComputeRecords.size(), 0,localActivePatches.size(), simParams->alpha_cutoff-simParams->fsMax, simParams->coulomb_radius_offset, lata, latb, latc, stream2 ); } //cuda_load_forces(forces, (doSlow ? slow_forces : 0 ), // 0,num_local_atom_records); //cuda_load_virials(virials, doSlow); // slow_virials follows virials cudaEventRecord(end_local_download, stream2); if ( workStarted == 2 ) { GBISP("C.N.CUDA[%d]::recvYieldDeviceL: adding POLL \ cuda_check_local_progress\n", CkMyPe()) CUDA_POLL(cuda_check_local_progress,this); } if ( shared_gpu && ! mergegrids ) { GBISP("C.N.CUDA[%d]::recvYieldDeviceL: adding POLL \ cuda_check_local_calc\n", CkMyPe()) CUDA_POLL(cuda_check_local_calc,this); break; } } // !GBIS or gbisPhase==1 if (simParams->GBISOn) { if (gbisPhase == 1) { //GBIS P1 Kernel launched in previous code block } else if (gbisPhase == 2) { GBISP("C.N.CUDA[%d]::recvYieldDeviceL: calling <<<P2>>>\n", CkMyPe()) cuda_GBIS_P2( 0,localComputeRecords.size(), 0,localActivePatches.size(), (simParams->alpha_cutoff-simParams->fsMax), simParams->cutoff, simParams->nonbondedScaling, simParams->kappa, (simParams->switchingActive ? simParams->switchingDist : -1.0), simParams->dielectric, simParams->solvent_dielectric, lata, latb, latc, doEnergy, doSlow, stream2 ); cudaEventRecord(end_local_download, stream2); if ( workStarted == 2 ) { CUDA_POLL(cuda_check_local_progress,this); } } else if (gbisPhase == 3) { GBISP("C.N.CUDA[%d]::recvYieldDeviceL: calling <<<P3>>>\n", CkMyPe()) if (doSlow) cuda_GBIS_P3( 0,localComputeRecords.size(), 0,localActivePatches.size(), (simParams->alpha_cutoff-simParams->fsMax), simParams->coulomb_radius_offset, simParams->nonbondedScaling, lata, latb, latc, stream2 ); cudaEventRecord(end_local_download, stream2); if ( workStarted == 2 ) { CUDA_POLL(cuda_check_local_progress,this); } } // phases } // GBISOn default: GBISP("C.N.CUDA[%d]::recvYieldDevice: case default\n", CkMyPe()) gpu_is_mine = 1; break; } // switch GBISP("C.N.CUDA[%d]::recvYieldDevice: DONE\n", CkMyPe()) }

void ComputeNonbondedCUDA::registerPatches (   )

Definition at line 814 of file ComputeNonbondedCUDA.C. References activePatches, ResizeArray< Elem >::add(), ResizeArray< Elem >::begin(), ComputeNonbondedCUDA::patch_record::bornRadBox, ProxyMgr::createProxy(), ComputeNonbondedCUDA::patch_record::dEdaSumBox, ComputeNonbondedCUDA::patch_record::dHdrPrefixBox, ComputeNonbondedCUDA::patch_record::forceBox, SimParameters::GBISOn, hostedPatches, ComputeNonbondedCUDA::patch_record::hostPe, ComputeNonbondedCUDA::patch_record::intRadBox, ComputeNonbondedCUDA::patch_record::isLocal, localHostedPatches, master, ProxyMgr::Object(), Node::Object(), ComputeNonbondedCUDA::patch_record::p, PatchMap::patch(), patchMap, patchRecords, ComputeNonbondedCUDA::patch_record::positionBox, ComputeNonbondedCUDA::patch_record::psiSumBox, Patch::registerBornRadPickup(), Patch::registerDEdaSumDeposit(), Patch::registerDHdrPrefixPickup(), Patch::registerForceDeposit(), Patch::registerIntRadPickup(), Patch::registerPositionPickup(), Patch::registerPsiSumDeposit(), remoteHostedPatches, Compute::setNumPatches(), Node::simParameters, simParams, and ResizeArray< Elem >::size(). Referenced by assignPatches(), and ComputeNonbondedCUDA(). { SimParameters *simParams = Node::Object()->simParameters; int npatches = master->activePatches.size(); int *pids = master->activePatches.begin(); patch_record *recs = master->patchRecords.begin(); for ( int i=0; i<npatches; ++i ) { int pid = pids[i]; patch_record &pr = recs[pid]; if ( pr.hostPe == CkMyPe() ) { //if ( ! reduction ) { // reduction = ReductionMgr::Object()->willSubmit(REDUCTIONS_BASIC); //} hostedPatches.add(pid); if ( pr.isLocal ) { localHostedPatches.add(pid); } else { remoteHostedPatches.add(pid); } ProxyMgr::Object()->createProxy(pid); pr.p = patchMap->patch(pid); pr.positionBox = pr.p->registerPositionPickup(this); pr.forceBox = pr.p->registerForceDeposit(this); if (simParams->GBISOn) { pr.intRadBox = pr.p->registerIntRadPickup(this); pr.psiSumBox = pr.p->registerPsiSumDeposit(this); pr.bornRadBox = pr.p->registerBornRadPickup(this); pr.dEdaSumBox = pr.p->registerDEdaSumDeposit(this); pr.dHdrPrefixBox = pr.p->registerDHdrPrefixPickup(this); } } } if ( master == this ) setNumPatches(activePatches.size()); else setNumPatches(hostedPatches.size()); CkPrintf("Pe %d hosts %d local and %d remote patches for pe %d\n", CkMyPe(), localHostedPatches.size(), remoteHostedPatches.size(), masterPe); }

void ComputeNonbondedCUDA::requirePatch (  int  pid  )

Definition at line 775 of file ComputeNonbondedCUDA.C. References activePatches, ResizeArray< Elem >::add(), ComputeNonbondedCUDA::patch_record::bornRad, computesChanged, ComputeNonbondedCUDA::patch_record::dEdaSum, ComputeNonbondedCUDA::patch_record::dHdrPrefix, ComputeNonbondedCUDA::patch_record::f, ComputeNonbondedCUDA::patch_record::hostPe, PatchMap::index_a(), PatchMap::index_b(), PatchMap::index_c(), ComputeNonbondedCUDA::patch_record::intRad, ComputeNonbondedCUDA::patch_record::isLocal, ResizeArray< Elem >::item(), localActivePatches, PatchMap::node(), ComputeNonbondedCUDA::patch_record::patchID, patchMap, patchRecords, ComputeNonbondedCUDA::patch_record::psiSum, ComputeNonbondedCUDA::patch_record::r, ComputeNonbondedCUDA::patch_record::refCount, remoteActivePatches, ComputeNonbondedCUDA::patch_record::x, and ComputeNonbondedCUDA::patch_record::xExt. Referenced by register_cuda_compute_pair(), and register_cuda_compute_self(). { computesChanged = 1; patch_record &pr = patchRecords.item(pid); if ( pr.refCount == 0 ) { if ( mergegrids ) { pr.isLocal = 0; } else if ( CkNumNodes() < 2 ) { pr.isLocal = 1 & ( 1 ^ patchMap->index_a(pid) ^ patchMap->index_b(pid) ^ patchMap->index_c(pid) ); } else { pr.isLocal = ( CkNodeOf(patchMap->node(pid)) == CkMyNode() ); } if ( pr.isLocal ) { localActivePatches.add(pid); } else { remoteActivePatches.add(pid); } activePatches.add(pid); // setNumPatches(activePatches.size()); pr.patchID = pid; pr.hostPe = -1; // ProxyMgr::Object()->createProxy(pid); // pr.p = patchMap->patch(pid); // pr.positionBox = pr.p->registerPositionPickup(this); // pr.forceBox = pr.p->registerForceDeposit(this); pr.x = NULL; pr.xExt = NULL; pr.r = NULL; pr.f = NULL; pr.intRad = NULL; pr.psiSum = NULL; pr.bornRad = NULL; pr.dEdaSum = NULL; pr.dHdrPrefix = NULL; } pr.refCount += 1; }

---

Member Data Documentation

ResizeArray<int> ComputeNonbondedCUDA::activePatches

Definition at line 80 of file ComputeNonbondedCUDA.h. Referenced by assignPatches(), doWork(), registerPatches(), and requirePatch().

AtomMap* ComputeNonbondedCUDA::atomMap

Definition at line 97 of file ComputeNonbondedCUDA.h. Referenced by ComputeNonbondedCUDA().

ResizeArray<compute_record> ComputeNonbondedCUDA::computeRecords

Definition at line 83 of file ComputeNonbondedCUDA.h. Referenced by doWork().

GBReal* ComputeNonbondedCUDA::dEdaSumH

Definition at line 94 of file ComputeNonbondedCUDA.h. Referenced by doWork(), finishWork(), and recvYieldDevice().

int ComputeNonbondedCUDA::doEnergy

Definition at line 67 of file ComputeNonbondedCUDA.h. Referenced by noWork().

int ComputeNonbondedCUDA::doSlow

Definition at line 67 of file ComputeNonbondedCUDA.h. Referenced by noWork().

float4* ComputeNonbondedCUDA::forces

Definition at line 91 of file ComputeNonbondedCUDA.h. Referenced by doWork(), finishWork(), and recvYieldDevice().

ResizeArray<int> ComputeNonbondedCUDA::hostedPatches

Definition at line 81 of file ComputeNonbondedCUDA.h. Referenced by doWork(), noWork(), and registerPatches().

ComputeNonbondedCUDAKernel* ComputeNonbondedCUDA::kernel

Definition at line 100 of file ComputeNonbondedCUDA.h.

Lattice ComputeNonbondedCUDA::lattice

Definition at line 66 of file ComputeNonbondedCUDA.h. Referenced by doWork(), noWork(), and recvYieldDevice().

ResizeArray<int> ComputeNonbondedCUDA::localActivePatches

Definition at line 80 of file ComputeNonbondedCUDA.h. Referenced by doWork(), recvYieldDevice(), and requirePatch().

ResizeArray<compute_record> ComputeNonbondedCUDA::localComputeRecords

Definition at line 84 of file ComputeNonbondedCUDA.h. Referenced by doWork(), recvYieldDevice(), register_cuda_compute_pair(), and register_cuda_compute_self().

ResizeArray<int> ComputeNonbondedCUDA::localHostedPatches

Definition at line 81 of file ComputeNonbondedCUDA.h. Referenced by finishWork(), registerPatches(), and ComputeMgr::sendNonbondedCUDASlaveEnqueue().

LocalWorkMsg* ComputeNonbondedCUDA::localWorkMsg2

Definition at line 63 of file ComputeNonbondedCUDA.h. Referenced by ComputeNonbondedCUDA(), and ComputeMgr::sendNonbondedCUDASlaveEnqueue().

ComputeNonbondedCUDA* ComputeNonbondedCUDA::master

Definition at line 102 of file ComputeNonbondedCUDA.h. Referenced by ComputeNonbondedCUDA(), doWork(), finishWork(), noWork(), and registerPatches().

int ComputeNonbondedCUDA::masterPe

Definition at line 103 of file ComputeNonbondedCUDA.h. Referenced by ComputeNonbondedCUDA(), and noWork().

int ComputeNonbondedCUDA::num_atom_records

Definition at line 86 of file ComputeNonbondedCUDA.h. Referenced by doWork().

int ComputeNonbondedCUDA::num_force_records

Definition at line 89 of file ComputeNonbondedCUDA.h. Referenced by doWork().

int ComputeNonbondedCUDA::num_local_atom_records

Definition at line 87 of file ComputeNonbondedCUDA.h. Referenced by doWork().

int ComputeNonbondedCUDA::num_remote_atom_records

Definition at line 88 of file ComputeNonbondedCUDA.h. Referenced by doWork().

int ComputeNonbondedCUDA::numSlaves

Definition at line 107 of file ComputeNonbondedCUDA.h. Referenced by assignPatches().

PatchMap* ComputeNonbondedCUDA::patchMap

Definition at line 96 of file ComputeNonbondedCUDA.h. Referenced by assignPatches(), ComputeNonbondedCUDA(), doWork(), register_cuda_compute_pair(), registerPatches(), and requirePatch().

ResizeArray<patch_record> ComputeNonbondedCUDA::patchRecords

Definition at line 82 of file ComputeNonbondedCUDA.h. Referenced by assignPatches(), doWork(), finishWork(), noWork(), register_cuda_compute_pair(), register_cuda_compute_self(), registerPatches(), and requirePatch().

GBReal* ComputeNonbondedCUDA::psiSumH

Definition at line 93 of file ComputeNonbondedCUDA.h. Referenced by doWork(), finishWork(), and recvYieldDevice().

SubmitReduction* ComputeNonbondedCUDA::reduction

Definition at line 98 of file ComputeNonbondedCUDA.h. Referenced by assignPatches(), ComputeNonbondedCUDA(), doWork(), finishWork(), and noWork().

ResizeArray<int> ComputeNonbondedCUDA::remoteActivePatches

Definition at line 80 of file ComputeNonbondedCUDA.h. Referenced by doWork(), recvYieldDevice(), and requirePatch().

ResizeArray<compute_record> ComputeNonbondedCUDA::remoteComputeRecords

Definition at line 84 of file ComputeNonbondedCUDA.h. Referenced by doWork(), recvYieldDevice(), register_cuda_compute_pair(), and register_cuda_compute_self().

ResizeArray<int> ComputeNonbondedCUDA::remoteHostedPatches

Definition at line 81 of file ComputeNonbondedCUDA.h. Referenced by finishWork(), and registerPatches().

int ComputeNonbondedCUDA::slaveIndex

Definition at line 104 of file ComputeNonbondedCUDA.h.

int* ComputeNonbondedCUDA::slavePes

Definition at line 106 of file ComputeNonbondedCUDA.h. Referenced by assignPatches(), ComputeNonbondedCUDA(), and messageFinishWork().

ComputeNonbondedCUDA** ComputeNonbondedCUDA::slaves

Definition at line 105 of file ComputeNonbondedCUDA.h. Referenced by assignPatches(), ComputeNonbondedCUDA(), and messageFinishWork().

float4* ComputeNonbondedCUDA::slow_forces

Definition at line 92 of file ComputeNonbondedCUDA.h. Referenced by doWork(), finishWork(), and recvYieldDevice().

int ComputeNonbondedCUDA::step

Definition at line 68 of file ComputeNonbondedCUDA.h. Referenced by finishWork(), and noWork().

int ComputeNonbondedCUDA::workStarted

Definition at line 65 of file ComputeNonbondedCUDA.h. Referenced by ComputeNonbondedCUDA(), doWork(), finishWork(), messageFinishWork(), noWork(), and recvYieldDevice().

---

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

• ComputeNonbondedCUDA.h
• ComputeNonbondedCUDA.C

---