CudaComputeNonbondedInteractions.h File Reference

#include "CudaUtils.h"

Go to the source code of this file.

Functions
template<bool doEnergy>
__device__ __forceinline__ void	cudaNBForce_Vdw_EnergySwitch (const float r2, const float rinv6, const float rinv8, const float2 ljab, const CudaNBConstants c, float &f_vdw, float &energyVdw)
template<bool doEnergy>
__device__ __forceinline__ void	cudaNBForce_PMESlowAndFast_C1 (const float r2, const float rinv, const float rinv2, const float rinv3, const float charge, const CudaNBConstants c, float &f_elec, float &energyElec, float &energySlow)
template<bool doEnergy>
__device__ __forceinline__ void	cudaNBForce_PMEFast_C1 (const float r2, const float rinv, const float rinv2, const float rinv3, const float charge, const CudaNBConstants c, float &f_elec, float &energyElec)
template<bool doEnergy>
__device__ __forceinline__ void	cudaNBForce_PMESlow_C1 (const float r2, const float rinv, const float rinv2, const float rinv3, const float charge, const CudaNBConstants c, float &fSlow, float &energySlow)
template<bool doEnergy>
__device__ __forceinline__ void	cudaModExclForce_PMESlow_C1 (const float r2, const float rinv, const float rinv2, const float rinv3, const float charge, const CudaNBConstants c, float &f_elec, float &energySlow)
template<bool doEnergy, bool doSlow>
__device__ __forceinline__ void	cudaNBForceMagCalc_VdwEnergySwitch_PMEC1 (const float r2, const float rinv, const float charge, const float2 ljab, const CudaNBConstants c, float &f, float &fSlow, float &energyVdw, float &energyElec, float &energySlow)
template<bool doEnergy>
__device__ __forceinline__ void	cudaModExclForceMagCalc_VdwEnergySwitch_PMEC1 (const int doSlow, const int doElec, const float r2, const float rinv, const float charge, const float2 ljab, const CudaNBConstants c, float &f, float &fSlow, float &energyVdw, float &energyElec, float &energySlow)

Function Documentation

cudaModExclForce_PMESlow_C1()

template<bool doEnergy>

__device__ __forceinline__ void cudaModExclForce_PMESlow_C1	(	const float	r2,
		const float	rinv,
		const float	rinv2,
		const float	rinv3,
		const float	charge,
		const CudaNBConstants	c,
		float &	f_elec,
		float &	energySlow
	)

Definition at line 170 of file CudaComputeNonbondedInteractions.h.

References CudaNBConstants::e_0, CudaNBConstants::e_1, and CudaNBConstants::e_2.

  {
  const float elec_slow = -1.0f * c.e_0;
  f_elec = charge * elec_slow;

  if (doEnergy) {
    float slow_energy = 0.5f * c.e_2 * (3.0f - r2 * c.e_1);
    energySlow += charge * slow_energy;
  }
}

cudaModExclForceMagCalc_VdwEnergySwitch_PMEC1()

template<bool doEnergy>

__device__ __forceinline__ void cudaModExclForceMagCalc_VdwEnergySwitch_PMEC1	(	const int	doSlow,
		const int	doElec,
		const float	r2,
		const float	rinv,
		const float	charge,
		const float2	ljab,
		const CudaNBConstants	c,
		float &	f,
		float &	fSlow,
		float &	energyVdw,
		float &	energyElec,
		float &	energySlow
	)

Definition at line 223 of file CudaComputeNonbondedInteractions.h.

                                                          {

  const float rinv2 = rinv * rinv;
  const float rinv3 = rinv * rinv2;
  const float rinv6 = rinv3 * rinv3;
  const float rinv8 = rinv6 * rinv2;

  float f_vdw;
  cudaNBForce_Vdw_EnergySwitch<doEnergy>(r2, rinv6, rinv8, ljab, c, f_vdw, energyVdw);
  // Sign corrections. The force tables are flipped and exclusions kernel uses xyz.i - xyz.j while
  // the nonbonded kernel uses xyz.j - xyz.i
  f = -1.0f * f_vdw;

  // Electrostatics
  if (doElec) {
    float f_elec;
    cudaNBForce_PMEFast_C1<doEnergy>(r2, rinv, rinv2, rinv3, charge, c, f_elec, energyElec);

    f += f_elec;
    energyElec *= -1.0f;

    if (doSlow) {
      cudaModExclForce_PMESlow_C1<doEnergy>(r2, rinv, rinv2, rinv3, charge, c, fSlow, energySlow);
      energySlow *= -1.0f;
    }
  }
}  

cudaNBForce_PMEFast_C1()

template<bool doEnergy>

__device__ __forceinline__ void cudaNBForce_PMEFast_C1	(	const float	r2,
		const float	rinv,
		const float	rinv2,
		const float	rinv3,
		const float	charge,
		const CudaNBConstants	c,
		float &	f_elec,
		float &	energyElec
	)

Definition at line 109 of file CudaComputeNonbondedInteractions.h.

References CudaNBConstants::e_0, CudaNBConstants::e_1, and CudaNBConstants::e_2.

  {
  const float elec_fast = -1.0f * rinv3 + c.e_0;
  f_elec = charge * elec_fast;

  if (doEnergy) {
    float slow_energy = 0.5f * c.e_2 * (3.0f - r2 * c.e_1);
    float fast_energy = rinv - slow_energy;
    energyElec += charge * fast_energy;
  }
}

cudaNBForce_PMESlow_C1()

template<bool doEnergy>

__device__ __forceinline__ void cudaNBForce_PMESlow_C1	(	const float	r2,
		const float	rinv,
		const float	rinv2,
		const float	rinv3,
		const float	charge,
		const CudaNBConstants	c,
		float &	fSlow,
		float &	energySlow
	)

Definition at line 133 of file CudaComputeNonbondedInteractions.h.

References CudaNBConstants::e_0, CudaNBConstants::e_1, CudaNBConstants::e_2, CudaNBConstants::ewald_0, and CudaNBConstants::ewald_1.

  {
  fSlow = charge;

  const float r = sqrtf(r2);
  const float elec_a = r * c.ewald_0;
  // very expensive
  const float elec_b = erfcf(elec_a);
  const float elec_exp = expf(-1.0f * elec_a * elec_a);
  const float corr_grad = elec_b + c.ewald_1 * r * elec_exp;

  const float elec_slow = -1.0f * c.e_0 * r;
  const float scor_grad = (elec_slow - (corr_grad - 1.0f) * rinv2)*rinv;

  if (doEnergy) {
    float slow_energy = 0.5f * c.e_2 * (3.0f - r2 * c.e_1);
    const float corr_energy = elec_b;
    const float scor_energy = slow_energy + (corr_energy - 1.0f) * rinv;
    energySlow += fSlow * scor_energy;
  }

  fSlow *= scor_grad;
}

cudaNBForce_PMESlowAndFast_C1()

template<bool doEnergy>

__device__ __forceinline__ void cudaNBForce_PMESlowAndFast_C1	(	const float	r2,
		const float	rinv,
		const float	rinv2,
		const float	rinv3,
		const float	charge,
		const CudaNBConstants	c,
		float &	f_elec,
		float &	energyElec,
		float &	energySlow
	)

Definition at line 73 of file CudaComputeNonbondedInteractions.h.

References CudaNBConstants::e_0, CudaNBConstants::e_1, CudaNBConstants::e_2, CudaNBConstants::ewald_0, CudaNBConstants::ewald_1, and CudaNBConstants::slowScale.

  {
  const float elec_fast = -1.0f * rinv3 + c.e_0;
  const float r = sqrtf(r2);
  const float elec_a = r * c.ewald_0;
  const float elec_exp = expf(-1.0f * elec_a * elec_a);
  const float elec_b = erfcf(elec_a);
  const float corr_grad = elec_b + c.ewald_1 * r * elec_exp;

  const float elec_slow = -1.0f * c.e_0 * r;
  const float scor_grad = (elec_slow - (corr_grad - 1.0f) * rinv2)*rinv;

  if (doEnergy) {
    float slow_energy = 0.5f * c.e_2 * (3.0f - r2 * c.e_1);
    float fast_energy = rinv - slow_energy;
    energyElec += charge * fast_energy;
    const float corr_energy = elec_b;
    const float scor_energy = slow_energy + (corr_energy - 1.0f) * rinv;
    energySlow += charge * scor_energy;
  }
  f_elec = charge * (elec_fast + scor_grad * c.slowScale);
}

cudaNBForce_Vdw_EnergySwitch()

template<bool doEnergy>

__device__ __forceinline__ void cudaNBForce_Vdw_EnergySwitch	(	const float	r2,
		const float	rinv6,
		const float	rinv8,
		const float2	ljab,
		const CudaNBConstants	c,
		float &	f_vdw,
		float &	energyVdw
	)

Definition at line 39 of file CudaComputeNonbondedInteractions.h.

References CudaNBConstants::lj_0, CudaNBConstants::lj_1, CudaNBConstants::lj_2, CudaNBConstants::lj_3, CudaNBConstants::lj_4, and CudaNBConstants::lj_5.

                                  {

  const float ab_r6 = ljab.x * rinv6 - ljab.y;
  const float w = ab_r6 * rinv6;
  const float dw_r = (ljab.x * rinv6 + ab_r6) * -6.0f * rinv8;

  float e_vdw;

  if (r2 > c.lj_5) {
    const float delta_r = (c.lj_4 - r2);
    const float s = delta_r * delta_r * (c.lj_0 + c.lj_1 * r2);
    const float ds_r = delta_r * (c.lj_3 + c.lj_2 * r2);
    f_vdw = w * ds_r + dw_r * s;
    if (doEnergy) e_vdw = w * s;
  } else {
    f_vdw = dw_r;
    if (doEnergy) e_vdw = w;
  }

  if (doEnergy) energyVdw += e_vdw;
}

cudaNBForceMagCalc_VdwEnergySwitch_PMEC1()

template<bool doEnergy, bool doSlow>

__device__ __forceinline__ void cudaNBForceMagCalc_VdwEnergySwitch_PMEC1	(	const float	r2,
		const float	rinv,
		const float	charge,
		const float2	ljab,
		const CudaNBConstants	c,
		float &	f,
		float &	fSlow,
		float &	energyVdw,
		float &	energyElec,
		float &	energySlow
	)

Definition at line 192 of file CudaComputeNonbondedInteractions.h.

                                                          {

  const float rinv2 = rinv * rinv;
  const float rinv3 = rinv * rinv2;
  const float rinv6 = rinv3 * rinv3;
  const float rinv8 = rinv6 * rinv2;

  float f_vdw;
  cudaNBForce_Vdw_EnergySwitch<doEnergy>(r2, rinv6, rinv8, ljab, c, f_vdw, energyVdw);

  // Electrostatics
  float f_elec;
  if (!doSlow) {
    cudaNBForce_PMEFast_C1<doEnergy>(r2, rinv, rinv2, rinv3, charge, c, f_elec, energyElec);
  } else {
    cudaNBForce_PMEFast_C1<doEnergy>(r2, rinv, rinv2, rinv3, charge, c, f_elec, energyElec);
    cudaNBForce_PMESlow_C1<doEnergy>(r2, rinv, rinv2, rinv3, charge, c, fSlow, energySlow);
  }
  f = f_elec + f_vdw;
}