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 125 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 584 of file CudaPmeSolverUtil.C.

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

585  :
586  PmeRealSpaceCompute(pmeGrid, jblock, kblock), deviceID(deviceID), stream(stream) {
587  if (dataSize < xsize*ysize*zsize)
588  NAMD_bug("CudaPmeRealSpaceCompute::CudaPmeRealSpaceCompute, insufficient dataSize");
589  cudaCheck(cudaSetDevice(deviceID));
590  d_atomsCapacity = 0;
591  d_atoms = NULL;
592  d_forceCapacity = 0;
593  d_force = NULL;
594  #ifdef NAMD_CUDA
595  tex_data = NULL;
596  tex_data_len = 0;
597  #else
598  grid_data = NULL;
599  grid_data_len = 0;
600  #endif
601  allocate_device<float>(&data, dataSize);
602  setupGridData(data, xsize*ysize*zsize);
603  cudaCheck(cudaEventCreate(&gatherForceEvent));
604 }
PmeRealSpaceCompute(PmeGrid pmeGrid, const int jblock, const int kblock)
__thread cudaStream_t stream
void NAMD_bug(const char *err_msg)
Definition: common.C:129
#define cudaCheck(stmt)
Definition: CudaUtils.h:95
CudaPmeRealSpaceCompute::~CudaPmeRealSpaceCompute ( )

Definition at line 609 of file CudaPmeSolverUtil.C.

References cudaCheck, and PmeRealSpaceCompute::data.

609  {
610  cudaCheck(cudaSetDevice(deviceID));
611  if (d_atoms != NULL) deallocate_device<CudaAtom>(&d_atoms);
612  if (d_force != NULL) deallocate_device<CudaForce>(&d_force);
613  // if (d_patches != NULL) deallocate_device<PatchInfo>(&d_patches);
614  // deallocate_device<double>(&d_selfEnergy);
615  deallocate_device<float>(&data);
616  cudaCheck(cudaEventDestroy(gatherForceEvent));
617 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:95

Member Function Documentation

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

Implements PmeRealSpaceCompute.

Definition at line 640 of file CudaPmeSolverUtil.C.

References atoms, cudaCheck, and PmeRealSpaceCompute::numAtoms.

640  {
641  cudaCheck(cudaSetDevice(deviceID));
642  this->numAtoms = numAtoms;
643 
644  // Reallocate device arrays as neccessary
645  reallocate_device<CudaAtom>(&d_atoms, &d_atomsCapacity, numAtoms, 1.5f);
646 
647  // Copy atom data to device
648  copy_HtoD<CudaAtom>(atoms, d_atoms, numAtoms, stream);
649 }
static __thread atom * atoms
__thread cudaStream_t stream
#define cudaCheck(stmt)
Definition: CudaUtils.h:95
void CudaPmeRealSpaceCompute::gatherForce ( Lattice lattice,
CudaForce force 
)
virtual

Implements PmeRealSpaceCompute.

Definition at line 751 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.

751  {
752  cudaCheck(cudaSetDevice(deviceID));
753 
754  // Re-allocate force array if needed
755  reallocate_device<CudaForce>(&d_force, &d_forceCapacity, numAtoms, 1.5f);
756 
757  gather_force((const float4*)d_atoms, numAtoms,
759  xsize, ysize, zsize, xsize, y0, z0, (pmeGrid.yBlocks == 1), (pmeGrid.zBlocks == 1),
760  data, pmeGrid.order, (float3*)d_force,
761 #ifdef NAMD_CUDA
762  gridTexObj,
763 #endif
764  stream);
765 
766  copy_DtoH<CudaForce>(d_force, force, numAtoms, stream);
767 
768  cudaCheck(cudaEventRecord(gatherForceEvent, stream));
769 }
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:95
void CudaPmeRealSpaceCompute::gatherForceSetCallback ( ComputePmeCUDADevice devicePtr_in)

Definition at line 703 of file CudaPmeSolverUtil.C.

References CcdCallBacksReset(), and cudaCheck.

703  {
704  cudaCheck(cudaSetDevice(deviceID));
705  devicePtr = devicePtr_in;
706  checkCount = 0;
707  CcdCallBacksReset(0, CmiWallTimer());
708  // Set the call back at 0.1ms
709  CcdCallFnAfter(cuda_gatherforce_check, this, 0.1);
710 }
void CcdCallBacksReset(void *ignored, double curWallTime)
#define cudaCheck(stmt)
Definition: CudaUtils.h:95
void CudaPmeRealSpaceCompute::spreadCharge ( Lattice lattice)
virtual

Implements PmeRealSpaceCompute.

Definition at line 654 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.

654  {
655  cudaCheck(cudaSetDevice(deviceID));
656 
657  // Clear grid
658  clear_device_array<float>(data, xsize*ysize*zsize, stream);
659 
660  spread_charge((const float4*)d_atoms, numAtoms,
661  pmeGrid.K1, pmeGrid.K2, pmeGrid.K3, xsize, ysize, zsize,
662  xsize, y0, z0, (pmeGrid.yBlocks == 1), (pmeGrid.zBlocks == 1),
663  data, pmeGrid.order, stream);
664 
665  // ncall++;
666 
667  // if (ncall == 1) writeRealToDisk(data, xsize*ysize*zsize, "data.txt");
668 }
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:95
void CudaPmeRealSpaceCompute::waitGatherForceDone ( )

Definition at line 712 of file CudaPmeSolverUtil.C.

References cudaCheck.

712  {
713  cudaCheck(cudaSetDevice(deviceID));
714  cudaCheck(cudaEventSynchronize(gatherForceEvent));
715 }
#define cudaCheck(stmt)
Definition: CudaUtils.h:95

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