Move pair potentials and change to traditional Functor

kernels/potentials/pair_potentials.cuh

@@ -0,0 +1,118 @@
+#ifndef POTENTIALS_CUH
+#define POTENTIALS_CUH
+
+#include "precision.hpp"
+#include "vec3.h"
+#include <cmath>
+#include <cstdio>
+#include <variant>
+
+#ifdef __CUDACC__
+#define CUDA_CALLABLE __host__ __device__
+#else
+#define CUDA_CALLABLE
+#endif
+
+/**
+ * Result struct for the Pair Potential
+ */
+struct ForceAndEnergy {
+  real energy;
+  Vec3<real> force;
+
+  CUDA_CALLABLE inline static ForceAndEnergy zero() {
+    return {0.0, {0.0, 0.0, 0.0}};
+  };
+};
+
+/**
+ * Calculate the Lennard-Jones energy and force for the current particle
+ * pair described by displacement vector r
+ */
+struct LennardJones {
+  real m_sigma;
+  real m_epsilon;
+  real m_rcutoffsq;
+
+  CUDA_CALLABLE LennardJones(real sigma, real epsilon, real rcutoff) {
+    m_sigma = sigma;
+    m_epsilon = epsilon;
+    m_rcutoffsq = rcutoff * rcutoff;
+  };
+
+  CUDA_CALLABLE ForceAndEnergy calc_force_and_energy(Vec3<real> r) {
+    real rmagsq = r.squared_norm2();
+    if (rmagsq < m_rcutoffsq && rmagsq > 0.0) {
+      real inv_rmag = 1 / sqrt(rmagsq);
+
+      // Pre-Compute the terms (doing this saves on multiple devisions/pow
+      // function call)
+      real sigma_r = m_sigma * inv_rmag;
+      real sigma_r6 = sigma_r * sigma_r * sigma_r * sigma_r * sigma_r * sigma_r;
+      real sigma_r12 = sigma_r6 * sigma_r6;
+
+      // Get the energy
+      real energy = 4.0 * m_epsilon * (sigma_r12 - sigma_r6);
+
+      // Get the force vector
+      real force_mag =
+          4.0 * m_epsilon *
+          (12.0 * sigma_r12 * inv_rmag - 6.0 * sigma_r6 * inv_rmag);
+      Vec3<real> force = r.scale(force_mag * inv_rmag);
+
+      return {energy, force};
+
+    } else {
+      return ForceAndEnergy::zero();
+    }
+  };
+};
+
+/**
+ * Calculate the Morse potential energy and force for the current particle pair
+ * described by displacement vector r
+ */
+struct Morse {
+  real m_D;         // Depth of the potential well
+  real m_a;         // Width of the potential
+  real m_r0;        // Equilibrium bond distance
+  real m_rcutoffsq; // Cutoff distance squared
+
+  CUDA_CALLABLE Morse(real D, real a, real r0, real rcutoff) {
+    m_D = D;
+    m_a = a;
+    m_r0 = r0;
+    m_rcutoffsq = rcutoff * rcutoff;
+  };
+
+  CUDA_CALLABLE ForceAndEnergy calc_force_and_energy(Vec3<real> r) {
+    real rmagsq = r.squared_norm2();
+    if (rmagsq < m_rcutoffsq && rmagsq > 0.0) {
+      real rmag = sqrt(rmagsq);
+      real dr = rmag - m_r0;
+
+      // Compute exponentials
+      real exp_a_dr = exp(-m_a * dr);
+      real exp_2a_dr = exp_a_dr * exp_a_dr;
+
+      // Energy: V(r) = D * (exp(-2a(r - r0)) - 2*exp(-a(r - r0)))
+      real energy = m_D * (exp_2a_dr - 2.0 * exp_a_dr);
+
+      // Force magnitude: F(r) = 2aD * (exp(-2a(r - r0)) - exp(-a(r - r0)))
+      real force_mag = 2.0 * m_a * m_D * (exp_2a_dr - exp_a_dr);
+
+      // Direction: normalized vector
+      Vec3<real> force = r.scale(force_mag / rmag);
+
+      return {energy, force};
+
+    } else {
+      return ForceAndEnergy::zero();
+    }
+  };
+};
+
+// Variant type for storing pair potential types
+using PairPotentials = std::variant<LennardJones, Morse>;
+
+#endif