CudaComputeGBISKernel Class Reference

#include <CudaComputeGBISKernel.h>

Public Member Functions
	CudaComputeGBISKernel (int deviceID)
	~CudaComputeGBISKernel ()
void	updateIntRad (const int atomStorageSize, float *intRad0H, float *intRadSH, cudaStream_t stream)
void	updateBornRad (const int atomStorageSize, float *bornRadH, cudaStream_t stream)
void	update_dHdrPrefix (const int atomStorageSize, float *dHdrPrefixH, cudaStream_t stream)
void	GBISphase1 (CudaTileListKernel &tlKernel, const int atomStorageSize, const float3 lata, const float3 latb, const float3 latc, const float a_cut, float *h_psiSum, cudaStream_t stream)
void	GBISphase2 (CudaTileListKernel &tlKernel, const int atomStorageSize, const bool doEnergy, const bool doSlow, const float3 lata, const float3 latb, const float3 latc, const float r_cut, const float scaling, const float kappa, const float smoothDist, const float epsilon_p, const float epsilon_s, float4 *d_forces, float *h_dEdaSum, cudaStream_t stream)
void	GBISphase3 (CudaTileListKernel &tlKernel, const int atomStorageSize, const float3 lata, const float3 latb, const float3 latc, const float a_cut, float4 *d_forces, cudaStream_t stream)

Detailed Description

Definition at line 5 of file CudaComputeGBISKernel.h.

Constructor & Destructor Documentation

CudaComputeGBISKernel::CudaComputeGBISKernel

(

int

deviceID

)

Definition at line 435 of file CudaComputeGBISKernel.cu.

References cudaCheck.

                                                         : deviceID(deviceID) {
  cudaCheck(cudaSetDevice(deviceID));

  intRad0 = NULL;
  intRad0Size = 0;

  intRadS = NULL;
  intRadSSize = 0;

  psiSum = NULL;
  psiSumSize = 0;

  bornRad = NULL;
  bornRadSize = 0;

  dEdaSum = NULL;
  dEdaSumSize = 0;

  dHdrPrefix = NULL;
  dHdrPrefixSize = 0;

}

CudaComputeGBISKernel::~CudaComputeGBISKernel

(

)

Definition at line 461 of file CudaComputeGBISKernel.cu.

References cudaCheck.

                                              {
  cudaCheck(cudaSetDevice(deviceID));
  if (intRad0 != NULL) deallocate_device<float>(&intRad0);
  if (intRadS != NULL) deallocate_device<float>(&intRadS);
  if (psiSum != NULL) deallocate_device<float>(&psiSum);
  if (bornRad != NULL) deallocate_device<float>(&bornRad);
  if (dEdaSum != NULL) deallocate_device<float>(&dEdaSum);
  if (dHdrPrefix != NULL) deallocate_device<float>(&dHdrPrefix);
}

Member Function Documentation

void CudaComputeGBISKernel::GBISphase1	(	CudaTileListKernel &	tlKernel,
		const int	atomStorageSize,
		const float3	lata,
		const float3	latb,
		const float3	latc,
		const float	a_cut,
		float *	h_psiSum,
		cudaStream_t	stream
	)

Definition at line 503 of file CudaComputeGBISKernel.cu.

References GBISParam< 1 >::a_cut, cudaCheck, cutoff2, deviceCUDA, FS_MAX, GBISKERNEL_NUM_WARP, CudaTileListKernel::get_xyzq(), DeviceCUDA::getMaxNumBlocks(), CudaTileListKernel::getNumTileListsGBIS(), CudaTileListKernel::getPatchPairs(), CudaTileListKernel::getTileJatomStartGBIS(), CudaTileListKernel::getTileListsGBIS(), lata, latb, latc, stream, and WARPSIZE.

                       {

  reallocate_device<float>(&psiSum, &psiSumSize, atomStorageSize, 1.2f);
  clear_device_array<float>(psiSum, atomStorageSize, stream);

  int nwarp = GBISKERNEL_NUM_WARP;
  int nthread = WARPSIZE*nwarp;
  int nblock = min(deviceCUDA->getMaxNumBlocks(), (tlKernel.getNumTileListsGBIS()-1)/nwarp+1);

  GBISParam<1> param;
  param.a_cut = a_cut;

  float cutoff2 = (a_cut + FS_MAX)*(a_cut + FS_MAX);
  GBIS_Kernel<false, false, 1> <<< nblock, nthread, 0, stream >>> (
    tlKernel.getNumTileListsGBIS(), tlKernel.getTileListsGBIS(), tlKernel.getTileJatomStartGBIS(),
    tlKernel.getPatchPairs(), lata, latb, latc, tlKernel.get_xyzq(), cutoff2,
    param, intRad0, intRadS, NULL, psiSum, NULL, NULL);

  cudaCheck(cudaGetLastError());

  copy_DtoH<float>(psiSum, h_psiSum, atomStorageSize, stream);
}

void CudaComputeGBISKernel::GBISphase2	(	CudaTileListKernel &	tlKernel,
		const int	atomStorageSize,
		const bool	doEnergy,
		const bool	doSlow,
		const float3	lata,
		const float3	latb,
		const float3	latc,
		const float	r_cut,
		const float	scaling,
		const float	kappa,
		const float	smoothDist,
		const float	epsilon_p,
		const float	epsilon_s,
		float4 *	d_forces,
		float *	h_dEdaSum,
		cudaStream_t	stream
	)

Definition at line 531 of file CudaComputeGBISKernel.cu.

References CALL, cudaCheck, deviceCUDA, GBISParam< 2 >::epsilon_p_i, GBISParam< 2 >::epsilon_s_i, GBISKERNEL_NUM_WARP, DeviceCUDA::getMaxNumBlocks(), CudaTileListKernel::getNumTileListsGBIS(), GBISParam< 2 >::kappa, GBISParam< 2 >::r_cut2, GBISParam< 2 >::r_cut_2, GBISParam< 2 >::r_cut_4, GBISParam< 2 >::scaling, CudaTileListKernel::setTileListVirialEnergyGBISLength(), GBISParam< 2 >::smoothDist, stream, and WARPSIZE.

                                                           {

  reallocate_device<float>(&dEdaSum, &dEdaSumSize, atomStorageSize, 1.2f);
  clear_device_array<float>(dEdaSum, atomStorageSize, stream);

  if (doEnergy) {
    tlKernel.setTileListVirialEnergyGBISLength(tlKernel.getNumTileListsGBIS());
  }

  int nwarp = GBISKERNEL_NUM_WARP;
  int nthread = WARPSIZE*nwarp;

  int nblock = min(deviceCUDA->getMaxNumBlocks(), (tlKernel.getNumTileListsGBIS()-1)/nwarp+1);

  GBISParam<2> param;
  param.r_cut2      = r_cut*r_cut;
  param.r_cut_2     = 1.f / param.r_cut2;
  param.r_cut_4     = 4.f*param.r_cut_2*param.r_cut_2;
  param.epsilon_s_i = 1.f / epsilon_s;
  param.epsilon_p_i = 1.f / epsilon_p;
  param.scaling     = scaling;
  param.kappa       = kappa;
  param.smoothDist  = smoothDist;

#define CALL(DOENERGY, DOSLOW) \
  GBIS_Kernel<DOENERGY, DOSLOW, 2> <<< nblock, nthread, 0, stream >>> (\
    tlKernel.getNumTileListsGBIS(), tlKernel.getTileListsGBIS(), tlKernel.getTileJatomStartGBIS(), \
      tlKernel.getPatchPairs(), lata, latb, latc, tlKernel.get_xyzq(), param.r_cut2, \
      param, bornRad, NULL, NULL, dEdaSum, d_forces, tlKernel.getTileListVirialEnergy())

  if (!doEnergy && !doSlow) CALL(false, false);
  if (!doEnergy &&  doSlow) CALL(false, true);
  if ( doEnergy && !doSlow) CALL(true,  false);
  if ( doEnergy &&  doSlow) CALL(true,  true);

  cudaCheck(cudaGetLastError());

  copy_DtoH<float>(dEdaSum, h_dEdaSum, atomStorageSize, stream);
}

void CudaComputeGBISKernel::GBISphase3	(	CudaTileListKernel &	tlKernel,
		const int	atomStorageSize,
		const float3	lata,
		const float3	latb,
		const float3	latc,
		const float	a_cut,
		float4 *	d_forces,
		cudaStream_t	stream
	)

Definition at line 579 of file CudaComputeGBISKernel.cu.

References GBISParam< 3 >::a_cut, cudaCheck, cutoff2, deviceCUDA, FS_MAX, GBISKERNEL_NUM_WARP, CudaTileListKernel::get_xyzq(), DeviceCUDA::getMaxNumBlocks(), CudaTileListKernel::getNumTileListsGBIS(), CudaTileListKernel::getPatchPairs(), CudaTileListKernel::getTileJatomStartGBIS(), CudaTileListKernel::getTileListsGBIS(), lata, latb, latc, and WARPSIZE.

                                         {

  int nwarp = GBISKERNEL_NUM_WARP;
  int nthread = WARPSIZE*nwarp;
  int nblock = min(deviceCUDA->getMaxNumBlocks(), (tlKernel.getNumTileListsGBIS()-1)/nwarp+1);

  GBISParam<3> param;
  param.a_cut = a_cut;

  float cutoff2 = (a_cut + FS_MAX)*(a_cut + FS_MAX);
  GBIS_Kernel<false, false, 3> <<< nblock, nthread, 0, stream >>> (
    tlKernel.getNumTileListsGBIS(), tlKernel.getTileListsGBIS(), tlKernel.getTileJatomStartGBIS(),
    tlKernel.getPatchPairs(), lata, latb, latc, tlKernel.get_xyzq(), cutoff2,
    param, intRad0, intRadS, dHdrPrefix, NULL, d_forces, NULL);

  cudaCheck(cudaGetLastError());
}

void CudaComputeGBISKernel::update_dHdrPrefix	(	const int	atomStorageSize,
		float *	dHdrPrefixH,
		cudaStream_t	stream
	)

Definition at line 495 of file CudaComputeGBISKernel.cu.

References dHdrPrefixH, and stream.

                                                                                                                {
  reallocate_device<float>(&dHdrPrefix, &dHdrPrefixSize, atomStorageSize, 1.2f);
  copy_HtoD<float>(dHdrPrefixH, dHdrPrefix, atomStorageSize, stream);
}

void CudaComputeGBISKernel::updateBornRad	(	const int	atomStorageSize,
		float *	bornRadH,
		cudaStream_t	stream
	)

Definition at line 487 of file CudaComputeGBISKernel.cu.

References bornRadH, and stream.

                                                                                                         {
  reallocate_device<float>(&bornRad, &bornRadSize, atomStorageSize, 1.2f);
  copy_HtoD<float>(bornRadH, bornRad, atomStorageSize, stream);
}

void CudaComputeGBISKernel::updateIntRad	(	const int	atomStorageSize,
		float *	intRad0H,
		float *	intRadSH,
		cudaStream_t	stream
	)

Definition at line 474 of file CudaComputeGBISKernel.cu.

References intRad0H, intRadSH, and stream.

                       {

  reallocate_device<float>(&intRad0, &intRad0Size, atomStorageSize, 1.2f);
  reallocate_device<float>(&intRadS, &intRadSSize, atomStorageSize, 1.2f);

  copy_HtoD<float>(intRad0H, intRad0, atomStorageSize, stream);
  copy_HtoD<float>(intRadSH, intRadS, atomStorageSize, stream);
}

---

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

• CudaComputeGBISKernel.h
• CudaComputeGBISKernel.cu