CudaPmeKSpaceCompute Class Reference

#include <CudaPmeSolverUtil.h>

Inheritance diagram for CudaPmeKSpaceCompute:
PmeKSpaceCompute

List of all members.

Classes

struct  EnergyVirial

Public Member Functions

 CudaPmeKSpaceCompute (PmeGrid pmeGrid, const int permutation, const int jblock, const int kblock, double kappa, int deviceID, cudaStream_t stream)
 ~CudaPmeKSpaceCompute ()
void solve (Lattice &lattice, const bool doEnergy, const bool doVirial, float *data)
double getEnergy ()
void getVirial (double *virial)
void energyAndVirialSetCallback (CudaPmePencilXYZ *pencilPtr)
void energyAndVirialSetCallback (CudaPmePencilZ *pencilPtr)

Detailed Description

Definition at line 59 of file CudaPmeSolverUtil.h.


Constructor & Destructor Documentation

CudaPmeKSpaceCompute::CudaPmeKSpaceCompute ( PmeGrid  pmeGrid,
const int  permutation,
const int  jblock,
const int  kblock,
double  kappa,
int  deviceID,
cudaStream_t  stream 
)

Definition at line 191 of file CudaPmeSolverUtil.C.

References PmeKSpaceCompute::bm1, PmeKSpaceCompute::bm2, PmeKSpaceCompute::bm3, cudaCheck, PmeGrid::K1, PmeGrid::K2, and PmeGrid::K3.

00192                                                                                        : 
00193   PmeKSpaceCompute(pmeGrid, permutation, jblock, kblock, kappa),
00194   deviceID(deviceID), stream(stream) {
00195 
00196   cudaCheck(cudaSetDevice(deviceID));
00197 
00198   // Copy bm1 -> prefac_x on GPU memory
00199   float *bm1f = new float[pmeGrid.K1];
00200   float *bm2f = new float[pmeGrid.K2];
00201   float *bm3f = new float[pmeGrid.K3];
00202   for (int i=0;i < pmeGrid.K1;i++) bm1f[i] = (float)bm1[i];
00203   for (int i=0;i < pmeGrid.K2;i++) bm2f[i] = (float)bm2[i];
00204   for (int i=0;i < pmeGrid.K3;i++) bm3f[i] = (float)bm3[i];
00205   allocate_device<float>(&d_bm1, pmeGrid.K1);
00206   allocate_device<float>(&d_bm2, pmeGrid.K2);
00207   allocate_device<float>(&d_bm3, pmeGrid.K3);
00208   copy_HtoD_sync<float>(bm1f, d_bm1, pmeGrid.K1);
00209   copy_HtoD_sync<float>(bm2f, d_bm2, pmeGrid.K2);
00210   copy_HtoD_sync<float>(bm3f, d_bm3, pmeGrid.K3);
00211   delete [] bm1f;
00212   delete [] bm2f;
00213   delete [] bm3f;
00214   allocate_device<EnergyVirial>(&d_energyVirial, 1);
00215   allocate_host<EnergyVirial>(&h_energyVirial, 1);
00216   // cudaCheck(cudaEventCreateWithFlags(&copyEnergyVirialEvent, cudaEventDisableTiming));
00217   cudaCheck(cudaEventCreate(&copyEnergyVirialEvent));
00218   // ncall = 0;
00219 }

CudaPmeKSpaceCompute::~CudaPmeKSpaceCompute (  ) 

Definition at line 221 of file CudaPmeSolverUtil.C.

References cudaCheck.

00221                                             {
00222   cudaCheck(cudaSetDevice(deviceID));
00223   deallocate_device<float>(&d_bm1);
00224   deallocate_device<float>(&d_bm2);
00225   deallocate_device<float>(&d_bm3);
00226   deallocate_device<EnergyVirial>(&d_energyVirial);
00227   deallocate_host<EnergyVirial>(&h_energyVirial);
00228   cudaCheck(cudaEventDestroy(copyEnergyVirialEvent));
00229 }


Member Function Documentation

void CudaPmeKSpaceCompute::energyAndVirialSetCallback ( CudaPmePencilZ pencilPtr  ) 

Definition at line 391 of file CudaPmeSolverUtil.C.

References CcdCallBacksReset(), and cudaCheck.

00391                                                                                {
00392   cudaCheck(cudaSetDevice(deviceID));
00393   pencilXYZPtr = NULL;
00394   pencilZPtr = pencilPtr;
00395   checkCount = 0;
00396   CcdCallBacksReset(0, CmiWallTimer());
00397   // Set the call back at 0.1ms
00398   CcdCallFnAfter(energyAndVirialCheck, this, 0.1);
00399 }

void CudaPmeKSpaceCompute::energyAndVirialSetCallback ( CudaPmePencilXYZ pencilPtr  ) 

Definition at line 381 of file CudaPmeSolverUtil.C.

References CcdCallBacksReset(), and cudaCheck.

00381                                                                                  {
00382   cudaCheck(cudaSetDevice(deviceID));
00383   pencilXYZPtr = pencilPtr;
00384   pencilZPtr = NULL;
00385   checkCount = 0;
00386   CcdCallBacksReset(0, CmiWallTimer());
00387   // Set the call back at 0.1ms
00388   CcdCallFnAfter(energyAndVirialCheck, this, 0.1);
00389 }

double CudaPmeKSpaceCompute::getEnergy (  )  [virtual]

Implements PmeKSpaceCompute.

Definition at line 401 of file CudaPmeSolverUtil.C.

00401                                        {
00402   return h_energyVirial->energy;
00403 }

void CudaPmeKSpaceCompute::getVirial ( double *  virial  )  [virtual]

Implements PmeKSpaceCompute.

Definition at line 405 of file CudaPmeSolverUtil.C.

References Perm_cX_Y_Z, Perm_Z_cX_Y, and PmeKSpaceCompute::permutation.

00405                                                    {
00406   if (permutation == Perm_Z_cX_Y) {
00407     // h_energyVirial->virial is storing ZZ, ZX, ZY, XX, XY, YY
00408     virial[0] = h_energyVirial->virial[3];
00409     virial[1] = h_energyVirial->virial[4];
00410     virial[2] = h_energyVirial->virial[1];
00411 
00412     virial[3] = h_energyVirial->virial[4];
00413     virial[4] = h_energyVirial->virial[5];
00414     virial[5] = h_energyVirial->virial[2];
00415 
00416     virial[6] = h_energyVirial->virial[1];
00417     virial[7] = h_energyVirial->virial[7];
00418     virial[8] = h_energyVirial->virial[0];
00419   } else if (permutation == Perm_cX_Y_Z) {
00420     // h_energyVirial->virial is storing XX, XY, XZ, YY, YZ, ZZ
00421     virial[0] = h_energyVirial->virial[0];
00422     virial[1] = h_energyVirial->virial[1];
00423     virial[2] = h_energyVirial->virial[2];
00424 
00425     virial[3] = h_energyVirial->virial[1];
00426     virial[4] = h_energyVirial->virial[3];
00427     virial[5] = h_energyVirial->virial[4];
00428 
00429     virial[6] = h_energyVirial->virial[2];
00430     virial[7] = h_energyVirial->virial[4];
00431     virial[8] = h_energyVirial->virial[5];
00432   }
00433 }

void CudaPmeKSpaceCompute::solve ( Lattice lattice,
const bool  doEnergy,
const bool  doVirial,
float *  data 
) [virtual]

Implements PmeKSpaceCompute.

Definition at line 231 of file CudaPmeSolverUtil.C.

References Lattice::a(), Lattice::a_r(), Lattice::b(), Lattice::b_r(), Lattice::c(), Lattice::c_r(), cudaCheck, PmeKSpaceCompute::j0, PmeKSpaceCompute::k0, PmeGrid::K1, PmeGrid::K2, PmeGrid::K3, PmeKSpaceCompute::kappa, NAMD_bug(), Perm_cX_Y_Z, Perm_Z_cX_Y, PmeKSpaceCompute::permutation, PmeKSpaceCompute::pmeGrid, scalar_sum(), PmeKSpaceCompute::size1, PmeKSpaceCompute::size2, PmeKSpaceCompute::size3, Lattice::volume(), Vector::x, Vector::y, and Vector::z.

00231                                                                                                         {
00232 #if 0
00233   // Check lattice to make sure it is updating for constant pressure
00234   fprintf(stderr, "K-SPACE LATTICE  %g %g %g  %g %g %g  %g %g %g\n",
00235       lattice.a().x, lattice.a().y, lattice.a().z,
00236       lattice.b().x, lattice.b().y, lattice.b().z,
00237       lattice.c().x, lattice.c().y, lattice.c().z);
00238 #endif
00239   cudaCheck(cudaSetDevice(deviceID));
00240 
00241   const bool doEnergyVirial = (doEnergy || doVirial);
00242 
00243   int nfft1, nfft2, nfft3;
00244   float *prefac1, *prefac2, *prefac3;
00245 
00246   BigReal volume = lattice.volume();
00247   Vector a_r = lattice.a_r();
00248   Vector b_r = lattice.b_r();
00249   Vector c_r = lattice.c_r();
00250   float recip1x, recip1y, recip1z;
00251   float recip2x, recip2y, recip2z;
00252   float recip3x, recip3y, recip3z;
00253 
00254   if (permutation == Perm_Z_cX_Y) {
00255     // Z, X, Y
00256     nfft1 = pmeGrid.K3;
00257     nfft2 = pmeGrid.K1;
00258     nfft3 = pmeGrid.K2;
00259     prefac1 = d_bm3;
00260     prefac2 = d_bm1;
00261     prefac3 = d_bm2;
00262     recip1x = c_r.z;
00263     recip1y = c_r.x;
00264     recip1z = c_r.y;
00265     recip2x = a_r.z;
00266     recip2y = a_r.x;
00267     recip2z = a_r.y;
00268     recip3x = b_r.z;
00269     recip3y = b_r.x;
00270     recip3z = b_r.y;
00271   } else if (permutation == Perm_cX_Y_Z) {
00272     // X, Y, Z
00273     nfft1 = pmeGrid.K1;
00274     nfft2 = pmeGrid.K2;
00275     nfft3 = pmeGrid.K3;
00276     prefac1 = d_bm1;
00277     prefac2 = d_bm2;
00278     prefac3 = d_bm3;
00279     recip1x = a_r.x;
00280     recip1y = a_r.y;
00281     recip1z = a_r.z;
00282     recip2x = b_r.x;
00283     recip2y = b_r.y;
00284     recip2z = b_r.z;
00285     recip3x = c_r.x;
00286     recip3y = c_r.y;
00287     recip3z = c_r.z;
00288   } else {
00289     NAMD_bug("CudaPmeKSpaceCompute::solve, invalid permutation");
00290   }
00291 
00292   // ncall++;
00293   // if (ncall == 1) {
00294   //   char filename[256];
00295   //   sprintf(filename,"dataf_%d_%d.txt",jblock,kblock);
00296   //   writeComplexToDisk((float2*)data, size1*size2*size3, filename, stream);
00297   // }
00298 
00299   // if (ncall == 1) {
00300   //   float2* h_data = new float2[size1*size2*size3];
00301   //   float2* d_data = (float2*)data;
00302   //   copy_DtoH<float2>(d_data, h_data, size1*size2*size3, stream);
00303   //   cudaCheck(cudaStreamSynchronize(stream));
00304   //   FILE *handle = fopen("dataf.txt", "w");
00305   //   for (int z=0;z < pmeGrid.K3;z++) {
00306   //     for (int y=0;y < pmeGrid.K2;y++) {
00307   //       for (int x=0;x < pmeGrid.K1/2+1;x++) {
00308   //         int i;
00309   //         if (permutation == Perm_cX_Y_Z) {
00310   //           i = x + y*size1 + z*size1*size2;
00311   //         } else {
00312   //           i = z + x*size1 + y*size1*size2;
00313   //         }
00314   //         fprintf(handle, "%f %f\n", h_data[i].x, h_data[i].y);
00315   //       }
00316   //     }
00317   //   }
00318   //   fclose(handle);
00319   //   delete [] h_data;
00320   // }
00321 
00322   // Clear energy and virial array if needed
00323   if (doEnergyVirial) clear_device_array<EnergyVirial>(d_energyVirial, 1, stream);
00324 
00325   scalar_sum(permutation == Perm_cX_Y_Z, nfft1, nfft2, nfft3, size1, size2, size3, kappa,
00326     recip1x, recip1y, recip1z, recip2x, recip2y, recip2z, recip3x, recip3y, recip3z,
00327     volume, prefac1, prefac2, prefac3, j0, k0, doEnergyVirial,
00328     &d_energyVirial->energy, d_energyVirial->virial, (float2*)data, 
00329     stream);
00330 
00331   // Copy energy and virial to host if needed
00332   if (doEnergyVirial) {
00333     copy_DtoH<EnergyVirial>(d_energyVirial, h_energyVirial, 1, stream);
00334     cudaCheck(cudaEventRecord(copyEnergyVirialEvent, stream));
00335     // cudaCheck(cudaStreamSynchronize(stream));
00336   }
00337 
00338 }


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

Generated on 24 May 2020 for NAMD by  doxygen 1.6.1