NAMD
Public Member Functions | List of all members
CudaPmeRealSpaceCompute Class Reference

#include <CudaPmeSolverUtil.h>

Inheritance diagram for CudaPmeRealSpaceCompute:
PmeRealSpaceCompute

Public Member Functions

 CudaPmeRealSpaceCompute (PmeGrid pmeGrid, const int jblock, const int kblock, int deviceID, cudaStream_t stream)
 
 ~CudaPmeRealSpaceCompute ()
 
void copyAtoms (const int numAtoms, const CudaAtom *atoms)
 
void spreadCharge (Lattice &lattice)
 
void gatherForce (Lattice &lattice, CudaForce *force)
 
void gatherForceSetCallback (ComputePmeCUDADevice *devicePtr_in)
 
void waitGatherForceDone ()
 
- Public Member Functions inherited from PmeRealSpaceCompute
 PmeRealSpaceCompute (PmeGrid pmeGrid, const int jblock, const int kblock)
 
virtual ~PmeRealSpaceCompute ()
 
float * getData ()
 
int getDataSize ()
 

Additional Inherited Members

- Static Public Member Functions inherited from PmeRealSpaceCompute
static double calcGridCoord (const double x, const double recip11, const int nfftx)
 
static void calcGridCoord (const double x, const double y, const double z, const double recip11, const double recip22, const double recip33, const int nfftx, const int nffty, const int nfftz, double &frx, double &fry, double &frz)
 
static void calcGridCoord (const float x, const float y, const float z, const float recip11, const float recip22, const float recip33, const int nfftx, const int nffty, const int nfftz, float &frx, float &fry, float &frz)
 
static void calcGridCoord (const float x, const float y, const float z, const int nfftx, const int nffty, const int nfftz, float &frx, float &fry, float &frz)
 
static void calcGridCoord (const double x, const double y, const double z, const int nfftx, const int nffty, const int nfftz, double &frx, double &fry, double &frz)
 
- Protected Attributes inherited from PmeRealSpaceCompute
int numAtoms
 
PmeGrid pmeGrid
 
int y0
 
int z0
 
int xsize
 
int ysize
 
int zsize
 
int dataSize
 
float * data
 
const int jblock
 
const int kblock
 

Detailed Description

Definition at line 98 of file CudaPmeSolverUtil.h.

Constructor & Destructor Documentation

CudaPmeRealSpaceCompute::CudaPmeRealSpaceCompute ( PmeGrid  pmeGrid,
const int  jblock,
const int  kblock,
int  deviceID,
cudaStream_t  stream 
)

Definition at line 443 of file CudaPmeSolverUtil.C.

References cudaCheck, PmeRealSpaceCompute::data, PmeRealSpaceCompute::dataSize, NAMD_bug(), PmeRealSpaceCompute::xsize, PmeRealSpaceCompute::ysize, and PmeRealSpaceCompute::zsize.

444  :
445  PmeRealSpaceCompute(pmeGrid, jblock, kblock), deviceID(deviceID), stream(stream) {
446  if (dataSize < xsize*ysize*zsize)
447  NAMD_bug("CudaPmeRealSpaceCompute::CudaPmeRealSpaceCompute, insufficient dataSize");
448  cudaCheck(cudaSetDevice(deviceID));
449  d_atomsCapacity = 0;
450  d_atoms = NULL;
451  d_forceCapacity = 0;
452  d_force = NULL;
453  tex_data = NULL;
454  tex_data_len = 0;
455  allocate_device<float>(&data, dataSize);
456  setupGridTexture(data, xsize*ysize*zsize);
457  cudaCheck(cudaEventCreate(&gatherForceEvent));
458 }
PmeRealSpaceCompute(PmeGrid pmeGrid, const int jblock, const int kblock)
__thread cudaStream_t stream
void NAMD_bug(const char *err_msg)
Definition: common.C:123
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
CudaPmeRealSpaceCompute::~CudaPmeRealSpaceCompute ( )

Definition at line 463 of file CudaPmeSolverUtil.C.

References cudaCheck, and PmeRealSpaceCompute::data.

463  {
464  cudaCheck(cudaSetDevice(deviceID));
465  if (d_atoms != NULL) deallocate_device<CudaAtom>(&d_atoms);
466  if (d_force != NULL) deallocate_device<CudaForce>(&d_force);
467  // if (d_patches != NULL) deallocate_device<PatchInfo>(&d_patches);
468  // deallocate_device<double>(&d_selfEnergy);
469  deallocate_device<float>(&data);
470  cudaCheck(cudaEventDestroy(gatherForceEvent));
471 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79

Member Function Documentation

void CudaPmeRealSpaceCompute::copyAtoms ( const int  numAtoms,
const CudaAtom atoms 
)
virtual

Implements PmeRealSpaceCompute.

Definition at line 494 of file CudaPmeSolverUtil.C.

References atoms, cudaCheck, and PmeRealSpaceCompute::numAtoms.

494  {
495  cudaCheck(cudaSetDevice(deviceID));
496  this->numAtoms = numAtoms;
497 
498  // Reallocate device arrays as neccessary
499  reallocate_device<CudaAtom>(&d_atoms, &d_atomsCapacity, numAtoms, 1.5f);
500 
501  // Copy atom data to device
502  copy_HtoD<CudaAtom>(atoms, d_atoms, numAtoms, stream);
503 }
static __thread atom * atoms
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
void CudaPmeRealSpaceCompute::gatherForce ( Lattice lattice,
CudaForce force 
)
virtual

Implements PmeRealSpaceCompute.

Definition at line 589 of file CudaPmeSolverUtil.C.

References cudaCheck, gather_force(), PmeGrid::K1, PmeGrid::K2, PmeGrid::K3, PmeRealSpaceCompute::numAtoms, PmeGrid::order, PmeRealSpaceCompute::pmeGrid, PmeRealSpaceCompute::xsize, PmeRealSpaceCompute::y0, PmeGrid::yBlocks, PmeRealSpaceCompute::ysize, PmeRealSpaceCompute::z0, PmeGrid::zBlocks, and PmeRealSpaceCompute::zsize.

589  {
590  cudaCheck(cudaSetDevice(deviceID));
591 
592  // Re-allocate force array if needed
593  reallocate_device<CudaForce>(&d_force, &d_forceCapacity, numAtoms, 1.5f);
594 
595  gather_force((const float4*)d_atoms, numAtoms,
597  xsize, ysize, zsize, xsize, y0, z0, (pmeGrid.yBlocks == 1), (pmeGrid.zBlocks == 1),
598  data, pmeGrid.order, (float3*)d_force,
599  gridTexObj,
600  stream);
601 
602  copy_DtoH<CudaForce>(d_force, force, numAtoms, stream);
603 
604  cudaCheck(cudaEventRecord(gatherForceEvent, stream));
605 }
int zBlocks
Definition: PmeBase.h:22
int K2
Definition: PmeBase.h:18
int K1
Definition: PmeBase.h:18
__thread cudaStream_t stream
int yBlocks
Definition: PmeBase.h:22
int order
Definition: PmeBase.h:20
__global__ void gather_force(const float4 *xyzq, const int ncoord, const int nfftx, const int nffty, const int nfftz, const int xsize, const int ysize, const int zsize, const int xdim, const int y00, const int z00, const float *data, const cudaTextureObject_t gridTexObj, const int stride, FT *force)
int K3
Definition: PmeBase.h:18
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
void CudaPmeRealSpaceCompute::gatherForceSetCallback ( ComputePmeCUDADevice devicePtr_in)

Definition at line 557 of file CudaPmeSolverUtil.C.

References CcdCallBacksReset(), and cudaCheck.

557  {
558  cudaCheck(cudaSetDevice(deviceID));
559  devicePtr = devicePtr_in;
560  checkCount = 0;
561  CcdCallBacksReset(0, CmiWallTimer());
562  // Set the call back at 0.1ms
563  CcdCallFnAfter(cuda_gatherforce_check, this, 0.1);
564 }
void CcdCallBacksReset(void *ignored, double curWallTime)
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
void CudaPmeRealSpaceCompute::spreadCharge ( Lattice lattice)
virtual

Implements PmeRealSpaceCompute.

Definition at line 508 of file CudaPmeSolverUtil.C.

References cudaCheck, PmeRealSpaceCompute::data, PmeGrid::K1, PmeGrid::K2, PmeGrid::K3, PmeRealSpaceCompute::numAtoms, PmeGrid::order, PmeRealSpaceCompute::pmeGrid, spread_charge(), PmeRealSpaceCompute::xsize, PmeRealSpaceCompute::y0, PmeGrid::yBlocks, PmeRealSpaceCompute::ysize, PmeRealSpaceCompute::z0, PmeGrid::zBlocks, and PmeRealSpaceCompute::zsize.

508  {
509  cudaCheck(cudaSetDevice(deviceID));
510 
511  // Clear grid
512  clear_device_array<float>(data, xsize*ysize*zsize, stream);
513 
514  spread_charge((const float4*)d_atoms, numAtoms,
515  pmeGrid.K1, pmeGrid.K2, pmeGrid.K3, xsize, ysize, zsize,
516  xsize, y0, z0, (pmeGrid.yBlocks == 1), (pmeGrid.zBlocks == 1),
517  data, pmeGrid.order, stream);
518 
519  // ncall++;
520 
521  // if (ncall == 1) writeRealToDisk(data, xsize*ysize*zsize, "data.txt");
522 }
int zBlocks
Definition: PmeBase.h:22
int K2
Definition: PmeBase.h:18
int K1
Definition: PmeBase.h:18
__thread cudaStream_t stream
int yBlocks
Definition: PmeBase.h:22
int order
Definition: PmeBase.h:20
int K3
Definition: PmeBase.h:18
void spread_charge(const float4 *atoms, const int numAtoms, const int nfftx, const int nffty, const int nfftz, const int xsize, const int ysize, const int zsize, const int xdim, const int y00, const int z00, const bool periodicY, const bool periodicZ, float *data, const int order, cudaStream_t stream)
#define cudaCheck(stmt)
Definition: CudaUtils.h:79
void CudaPmeRealSpaceCompute::waitGatherForceDone ( )

Definition at line 566 of file CudaPmeSolverUtil.C.

References cudaCheck.

566  {
567  cudaCheck(cudaSetDevice(deviceID));
568  cudaCheck(cudaEventSynchronize(gatherForceEvent));
569 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:79

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