CudaPmeKSpaceCompute Class Reference

#include <CudaPmeSolverUtil.h>

Inheritance diagram for CudaPmeKSpaceCompute:

Public Member Functions
	CudaPmeKSpaceCompute (PmeGrid pmeGrid, const int permutation, const int jblock, const int kblock, double kappa, int deviceID, cudaStream_t stream, unsigned int iGrid=0)
	~CudaPmeKSpaceCompute ()
void	solve (Lattice &lattice, const bool doEnergy, const bool doVirial, float *data)
void	waitEnergyAndVirial ()
double	getEnergy ()
void	getVirial (double *virial)
void	energyAndVirialSetCallback (CudaPmePencilXYZ *pencilPtr)
void	energyAndVirialSetCallback (CudaPmePencilZ *pencilPtr)
Public Member Functions inherited from PmeKSpaceCompute
	PmeKSpaceCompute (PmeGrid pmeGrid, const int permutation, const int jblock, const int kblock, double kappa, unsigned int multipleGridIndex=0)
virtual	~PmeKSpaceCompute ()
virtual void	setGrid (unsigned int iGrid)

Additional Inherited Members
Protected Attributes inherited from PmeKSpaceCompute
PmeGrid	pmeGrid
double *	bm1
double *	bm2
double *	bm3
double	kappa
const int	permutation
const int	jblock
const int	kblock
int	size1
int	size2
int	size3
int	j0
int	k0
unsigned int	multipleGridIndex

Detailed Description

Definition at line 73 of file CudaPmeSolverUtil.h.

Constructor & Destructor Documentation

CudaPmeKSpaceCompute()

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

Definition at line 236 of file CudaPmeSolverUtil.C.

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

                                                                                                           : 
  PmeKSpaceCompute(pmeGrid, permutation, jblock, kblock, kappa, iGrid),
  deviceID(deviceID), stream(stream) {

  cudaCheck(cudaSetDevice(deviceID));

  // Copy bm1 -> prefac_x on GPU memory
  float *bm1f = new float[pmeGrid.K1];
  float *bm2f = new float[pmeGrid.K2];
  float *bm3f = new float[pmeGrid.K3];
  for (int i=0;i < pmeGrid.K1;i++) bm1f[i] = (float)bm1[i];
  for (int i=0;i < pmeGrid.K2;i++) bm2f[i] = (float)bm2[i];
  for (int i=0;i < pmeGrid.K3;i++) bm3f[i] = (float)bm3[i];
  allocate_device<float>(&d_bm1, pmeGrid.K1);
  allocate_device<float>(&d_bm2, pmeGrid.K2);
  allocate_device<float>(&d_bm3, pmeGrid.K3);
  copy_HtoD_sync<float>(bm1f, d_bm1, pmeGrid.K1);
  copy_HtoD_sync<float>(bm2f, d_bm2, pmeGrid.K2);
  copy_HtoD_sync<float>(bm3f, d_bm3, pmeGrid.K3);
  delete [] bm1f;
  delete [] bm2f;
  delete [] bm3f;
  allocate_device<EnergyVirial>(&d_energyVirial, 1);
  allocate_host<EnergyVirial>(&h_energyVirial, 1);
  // cudaCheck(cudaEventCreateWithFlags(&copyEnergyVirialEvent, cudaEventDisableTiming));
  cudaCheck(cudaEventCreate(&copyEnergyVirialEvent));
  // ncall = 0;
}

~CudaPmeKSpaceCompute()

CudaPmeKSpaceCompute::~CudaPmeKSpaceCompute

(

)

Definition at line 266 of file CudaPmeSolverUtil.C.

References cudaCheck.

                                            {
  cudaCheck(cudaSetDevice(deviceID));
  deallocate_device<float>(&d_bm1);
  deallocate_device<float>(&d_bm2);
  deallocate_device<float>(&d_bm3);
  deallocate_device<EnergyVirial>(&d_energyVirial);
  deallocate_host<EnergyVirial>(&h_energyVirial);
  cudaCheck(cudaEventDestroy(copyEnergyVirialEvent));
}

Member Function Documentation

energyAndVirialSetCallback() [1/2]

void CudaPmeKSpaceCompute::energyAndVirialSetCallback

(

CudaPmePencilXYZ *

pencilPtr

)

Definition at line 477 of file CudaPmeSolverUtil.C.

References CcdCallBacksReset(), and cudaCheck.

                                                                                 {
  cudaCheck(cudaSetDevice(deviceID));
  pencilXYZPtr = pencilPtr;
  pencilZPtr = NULL;
  checkCount = 0;
  CcdCallBacksReset(0, CmiWallTimer());
  // Set the call back at 0.1ms
  CcdCallFnAfter(energyAndVirialCheck, this, 0.1);
}

energyAndVirialSetCallback() [2/2]

void CudaPmeKSpaceCompute::energyAndVirialSetCallback

(

CudaPmePencilZ *

pencilPtr

)

Definition at line 487 of file CudaPmeSolverUtil.C.

References CcdCallBacksReset(), and cudaCheck.

                                                                               {
  cudaCheck(cudaSetDevice(deviceID));
  pencilXYZPtr = NULL;
  pencilZPtr = pencilPtr;
  checkCount = 0;
  CcdCallBacksReset(0, CmiWallTimer());
  // Set the call back at 0.1ms
  CcdCallFnAfter(energyAndVirialCheck, this, 0.1);
}

getEnergy()

double CudaPmeKSpaceCompute::getEnergy ( )

virtual

Implements PmeKSpaceCompute.

Definition at line 497 of file CudaPmeSolverUtil.C.

                                       {
  return h_energyVirial->energy;
}

getVirial()

void CudaPmeKSpaceCompute::getVirial ( double *  virial )

virtual

Implements PmeKSpaceCompute.

Definition at line 501 of file CudaPmeSolverUtil.C.

References Perm_cX_Y_Z, Perm_Z_cX_Y, and PmeKSpaceCompute::permutation.

                                                   {
  if (permutation == Perm_Z_cX_Y) {
    // h_energyVirial->virial is storing ZZ, ZX, ZY, XX, XY, YY
    virial[0] = h_energyVirial->virial[3];
    virial[1] = h_energyVirial->virial[4];
    virial[2] = h_energyVirial->virial[1];

    virial[3] = h_energyVirial->virial[4];
    virial[4] = h_energyVirial->virial[5];
    virial[5] = h_energyVirial->virial[2];

    virial[6] = h_energyVirial->virial[1];
    virial[7] = h_energyVirial->virial[7];
    virial[8] = h_energyVirial->virial[0];
  } else if (permutation == Perm_cX_Y_Z) {
    // h_energyVirial->virial is storing XX, XY, XZ, YY, YZ, ZZ
    virial[0] = h_energyVirial->virial[0];
    virial[1] = h_energyVirial->virial[1];
    virial[2] = h_energyVirial->virial[2];

    virial[3] = h_energyVirial->virial[1];
    virial[4] = h_energyVirial->virial[3];
    virial[5] = h_energyVirial->virial[4];

    virial[6] = h_energyVirial->virial[2];
    virial[7] = h_energyVirial->virial[4];
    virial[8] = h_energyVirial->virial[5];
  }
#if 0
  fprintf(stderr, "AP PME VIRIAL =\n"
      "  %g  %g  %g\n  %g  %g  %g\n  %g %g %g\n",
      virial[0], virial[1], virial[2], virial[3], virial[4],
      virial[5], virial[6], virial[7], virial[8]);
#endif
}

solve()

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

virtual

Implements PmeKSpaceCompute.

Definition at line 276 of file CudaPmeSolverUtil.C.

References Lattice::a(), Lattice::a_r(), Lattice::b(), Lattice::b_r(), PmeKSpaceCompute::bm1, PmeKSpaceCompute::bm2, PmeKSpaceCompute::bm3, 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.

                                                                                                        {
#if 0
  // Check lattice to make sure it is updating for constant pressure
  fprintf(stderr, "K-SPACE LATTICE  %g %g %g  %g %g %g  %g %g %g\n",
      lattice.a().x, lattice.a().y, lattice.a().z,
      lattice.b().x, lattice.b().y, lattice.b().z,
      lattice.c().x, lattice.c().y, lattice.c().z);
#endif
  cudaCheck(cudaSetDevice(deviceID));

  const bool doEnergyVirial = (doEnergy || doVirial);

  int nfft1, nfft2, nfft3;
  float *prefac1, *prefac2, *prefac3;

  BigReal volume = lattice.volume();
  Vector a_r = lattice.a_r();
  Vector b_r = lattice.b_r();
  Vector c_r = lattice.c_r();
  float recip1x, recip1y, recip1z;
  float recip2x, recip2y, recip2z;
  float recip3x, recip3y, recip3z;

  if (permutation == Perm_Z_cX_Y) {
    // Z, X, Y
    nfft1 = pmeGrid.K3;
    nfft2 = pmeGrid.K1;
    nfft3 = pmeGrid.K2;
    prefac1 = d_bm3;
    prefac2 = d_bm1;
    prefac3 = d_bm2;
    recip1x = c_r.z;
    recip1y = c_r.x;
    recip1z = c_r.y;
    recip2x = a_r.z;
    recip2y = a_r.x;
    recip2z = a_r.y;
    recip3x = b_r.z;
    recip3y = b_r.x;
    recip3z = b_r.y;
  } else if (permutation == Perm_cX_Y_Z) {
    // X, Y, Z
    nfft1 = pmeGrid.K1;
    nfft2 = pmeGrid.K2;
    nfft3 = pmeGrid.K3;
    prefac1 = d_bm1;
    prefac2 = d_bm2;
    prefac3 = d_bm3;
    recip1x = a_r.x;
    recip1y = a_r.y;
    recip1z = a_r.z;
    recip2x = b_r.x;
    recip2y = b_r.y;
    recip2z = b_r.z;
    recip3x = c_r.x;
    recip3y = c_r.y;
    recip3z = c_r.z;
  } else {
    NAMD_bug("CudaPmeKSpaceCompute::solve, invalid permutation");
  }

  // ncall++;
  // if (ncall == 1) {
  //   char filename[256];
  //   sprintf(filename,"dataf_%d_%d.txt",jblock,kblock);
  //   writeComplexToDisk((float2*)data, size1*size2*size3, filename, stream);
  // }

  // if (ncall == 1) {
  //   float2* h_data = new float2[size1*size2*size3];
  //   float2* d_data = (float2*)data;
  //   copy_DtoH<float2>(d_data, h_data, size1*size2*size3, stream);
  //   cudaCheck(cudaStreamSynchronize(stream));
  //   FILE *handle = fopen("dataf.txt", "w");
  //   for (int z=0;z < pmeGrid.K3;z++) {
  //     for (int y=0;y < pmeGrid.K2;y++) {
  //       for (int x=0;x < pmeGrid.K1/2+1;x++) {
  //         int i;
  //         if (permutation == Perm_cX_Y_Z) {
  //           i = x + y*size1 + z*size1*size2;
  //         } else {
  //           i = z + x*size1 + y*size1*size2;
  //         }
  //         fprintf(handle, "%f %f\n", h_data[i].x, h_data[i].y);
  //       }
  //     }
  //   }
  //   fclose(handle);
  //   delete [] h_data;
  // }

  // Clear energy and virial array if needed
  if (doEnergyVirial) clear_device_array<EnergyVirial>(d_energyVirial, 1, stream);

#ifdef TESTPID
  if (1) {
    cudaCheck(cudaStreamSynchronize(stream));
    fprintf(stderr, "AP calling scalar sum\n");
    fprintf(stderr, "(permutation == Perm_cX_Y_Z) = %s\n",
        (permutation == Perm_cX_Y_Z ? "true" : "false"));
    fprintf(stderr, "nfft1=%d  nfft2=%d  nfft3=%d\n", nfft1, nfft2, nfft3);
    fprintf(stderr, "size1=%d  size2=%d  size3=%d\n", size1, size2, size3);
    fprintf(stderr, "kappa=%g\n", kappa);
    fprintf(stderr, "recip1x=%g  recip1y=%g  recip1z=%g\n",
        (double)recip1x, (double)recip1y, (double)recip1z);
    fprintf(stderr, "recip2x=%g  recip2y=%g  recip2z=%g\n",
        (double)recip2x, (double)recip2y, (double)recip2z);
    fprintf(stderr, "recip3x=%g  recip3y=%g  recip3z=%g\n",
        (double)recip3x, (double)recip3y, (double)recip3z);
    fprintf(stderr, "volume=%g\n", volume);
    fprintf(stderr, "j0=%d  k0=%d\n", j0, k0);
    float *bm1, *bm2, *bm3;
    allocate_host<float>(&bm1, nfft1);
    allocate_host<float>(&bm2, nfft2);
    allocate_host<float>(&bm3, nfft3);
    copy_DtoH<float>(prefac1, bm1, nfft1, stream);
    copy_DtoH<float>(prefac2, bm2, nfft2, stream);
    copy_DtoH<float>(prefac3, bm3, nfft3, stream);
    TestArray_write<float>("bm1_good.bin", "structure factor bm1 good",
        bm1, nfft1);
    TestArray_write<float>("bm2_good.bin", "structure factor bm2 good",
        bm2, nfft2);
    TestArray_write<float>("bm3_good.bin", "structure factor bm3 good",
        bm3, nfft3);
    deallocate_host<float>(&bm1);
    deallocate_host<float>(&bm2);
    deallocate_host<float>(&bm3);
  }
#endif

  scalar_sum(permutation == Perm_cX_Y_Z, nfft1, nfft2, nfft3, size1, size2, size3, kappa,
    recip1x, recip1y, recip1z, recip2x, recip2y, recip2z, recip3x, recip3y, recip3z,
    volume, prefac1, prefac2, prefac3, j0, k0, doEnergyVirial,
    &d_energyVirial->energy, d_energyVirial->virial, (float2*)data, 
    stream);
#ifdef TESTPID
  if (1) {
    cudaCheck(cudaStreamSynchronize(stream));
    fprintf(stderr, "AP SCALAR SUM\n");
    fprintf(stderr, "COPY DEVICE ARRAYS BACK TO HOST\n");
    int m = 2 * (nfft1/2 + 1) * nfft2 * nfft3;
    float *tran = 0;
    allocate_host<float>(&tran, m);
    copy_DtoH<float>((float*)data, tran, m, stream);
    cudaCheck(cudaStreamSynchronize(stream));
    TestArray_write<float>("tran_potential_grid_good.bin",
          "transformed potential grid good", tran, m);
    deallocate_host<float>(&tran);
  }
#endif

  // Copy energy and virial to host if needed
  if (doEnergyVirial) {
    copy_DtoH<EnergyVirial>(d_energyVirial, h_energyVirial, 1, stream);
    cudaCheck(cudaEventRecord(copyEnergyVirialEvent, stream));
    // cudaCheck(cudaStreamSynchronize(stream));
  }

}

waitEnergyAndVirial()

void CudaPmeKSpaceCompute::waitEnergyAndVirial

(

)

---

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

• CudaPmeSolverUtil.h
• CudaPmeSolverUtil.C