add cuda context

src/core/cuda/include/CudaAngleForce.cuh

@@ -1,9 +1,6 @@
-#ifndef CUDA_ANGLE_FORCE_H
-#define CUDA_ANGLE_FORCE_H
+#pragma once
 #include "system.h"
-__global__ void calc_angle_forces_kernel(int start, int end, angle_t *angles, coord_t *coords, cangle_t *cangles, dvel_t *dvelocities, double *energy_sum);
 
+void init_angle_force_kernel_data();
 double calc_angle_forces_host(int start, int end);
-
 void cleanup_angle_force();
-#endif

src/core/cuda/include/CudaBondForce.cuh

@@ -1,9 +1,6 @@
-#ifndef CUDA_BOND_FORCE_H
-#define CUDA_BOND_FORCE_H
+#pragma once
 #include "system.h"
 
-__global__ void calc_bond_forces_kernel(int start, int end, bond_t* bonds, coord_t *coords, cbond_t *cbonds, dvel_t *dvelocities, double *energy_sum);
+void init_bond_force_kernel_data();
 double calc_bond_forces_host(int start, int end);
-
 void cleanup_bond_force();
-#endif

src/core/cuda/include/CudaContext.cuh

@@ -0,0 +1,140 @@
+#pragma once
+
+#include <cuda_runtime.h>
+
+#include "system.h"
+#include "utils.h"
+
+class CudaContext {
+   public:
+    /*
+    Common data
+    */
+    coord_t* d_coords = nullptr;
+    dvel_t* d_dvelocities = nullptr;
+    vel_t* d_velocities = nullptr;
+
+    /*
+    Used in CudaAngleForce.cu
+    */
+    angle_t* d_angles = nullptr;
+    cangle_t* d_cangles = nullptr;
+
+    /*
+    Used in CudaBondForce.cu
+    */
+    bond_t* d_bonds = nullptr;
+    cbond_t* d_cbonds = nullptr;
+
+    /*
+    Used in CudaImproper2Force.cu
+    */
+    improper_t* d_impropers = nullptr;
+    cimproper_t* d_cimpropers = nullptr;
+
+    /*
+    Used in CudaLeapfrog.cu
+    */
+    atype_t* d_atypes = nullptr;
+    catype_t* d_catypes = nullptr;
+
+    /*
+    Used in CudaShakeConstraints.cu
+    */
+    int* d_mol_n_shakes = nullptr;
+    shake_bond_t* d_shake_bonds = nullptr;
+    double* d_winv = nullptr;
+    coord_t* d_xcoords = nullptr;
+
+    /*
+    Used in CudaNonbondedQQForce.cu
+    */
+    q_atom_t* d_q_atoms = nullptr;
+    q_charge_t* d_q_charges = nullptr;
+    int* d_LJ_matrix = nullptr;
+    bool* d_excluded = nullptr;
+    q_elscale_t* d_q_elscales = nullptr;
+    q_catype_t* d_q_catypes = nullptr;
+    q_atype_t* d_q_atypes = nullptr;
+    E_nonbonded_t* d_EQ_nonbond_qq = nullptr;
+    double* d_lambdas = nullptr;
+
+    /*
+    Used in CudaPolxWaterForce.cu
+    */
+    shell_t* d_wshells = nullptr;
+    /*
+    Used in CudaPshellForce.cu
+    */
+    bool* d_shell = nullptr;
+    coord_t* d_coords_top = nullptr;
+
+    /*
+    Used in CudaRestrangForce.cu
+    */
+    restrang_t* d_restrangs = nullptr;
+    E_restraint_t* d_EQ_restraint = nullptr;
+    restrdis_t* d_restrdists = nullptr;
+
+    /*
+    Used in CudaRestrposForce.cu
+    */
+    restrpos_t* d_restrspos = nullptr;
+
+    /*
+    Used in CudaRestrseqForce.cu
+    */
+    restrseq_t* d_restrseqs = nullptr;
+    bool* d_heavy = nullptr;
+
+    /*
+    Used in CudaRestrwallForce.cu
+    */
+    restrwall_t* d_restrwalls = nullptr;
+
+    /*
+    Used in CudaTorsionForce.cu
+    */
+    torsion_t* d_torsions = nullptr;
+    ctorsion_t* d_ctorsions = nullptr;
+
+    /*
+    Used in CudaNonbondedPPForce.cu
+    */
+    ccharge_t* d_ccharges;
+    charge_t* d_charges;
+    p_atom_t* d_p_atoms;
+
+    static CudaContext& instance() {
+        static CudaContext ctx;
+        return ctx;
+    }
+    void init();
+
+    template <typename T>
+    void sync_array_to_device(T* dst, const T* src, int count);
+
+    template <typename T>
+    void sync_array_to_host(T* dst, const T* src, int count);
+
+    void sync_all_to_device();
+    void sync_all_to_host();
+    void reset_energies();
+
+   private:
+    CudaContext() = default;
+
+    void free();
+
+    ~CudaContext() { free(); }
+    CudaContext(const CudaContext&) = delete;
+    CudaContext& operator=(const CudaContext&) = delete;
+};
+template <typename T>
+void CudaContext::sync_array_to_device(T* dst, const T* src, int count) {
+    check_cuda(cudaMemcpy(dst, src, count * sizeof(T), cudaMemcpyHostToDevice));
+}
+template <typename T>
+void CudaContext::sync_array_to_host(T* dst, const T* src, int count) {
+    check_cuda(cudaMemcpy(dst, src, count * sizeof(T), cudaMemcpyDeviceToHost));
+}

src/core/cuda/include/CudaImproper2Force.cuh

@@ -1,11 +1,6 @@
-#ifndef CUDA_IMPROPER2_FORCE
-#define CUDA_IMPROPER2_FORCE_H
+#pragma once
 #include "system.h"
 
-__global__ void calc_improper2_forces_kernel(int start, int end, improper_t* impropers, cimproper_t* cimpropers, coord_t* coords, dvel_t* dvelocities, double* energy_sum);
-
+void init_improper2_force_kernel_data();
 double calc_improper2_forces_host(int start, int end);
-
 void cleanup_improper2_force();
-
-#endif

src/core/cuda/include/CudaLeapfrog.cuh

@@ -1,8 +1,6 @@
-#ifndef CUDA_LEAPFROG_H
-#define CUDA_LEAPFROG_H
+#pragma once
 #include "system.h"
 
+void init_leapfrog_kernel_data();
 void calc_leapfrog_host();
 void cleanup_leapfrog();
-
-#endif

src/core/cuda/include/CudaNonbondedPPForce.cuh

@@ -0,0 +1,8 @@
+#pragma once
+#include "system.h"
+
+void init_nonbonded_pp_force_kernel_data();
+
+void calc_nonbonded_pp_forces_host_v2();
+
+void cleanup_nonbonded_pp_force();

src/core/cuda/include/CudaNonbondedPWForce.cuh

@@ -0,0 +1,7 @@
+#pragma once
+#include "system.h"
+
+void init_nonbonded_pw_force_kernel_data();
+void calc_nonbonded_pw_forces_host_v2();
+
+void cleanup_nonbonded_pw_force();

src/core/cuda/include/CudaNonbondedQPForce.cuh

@@ -0,0 +1,6 @@
+#pragma once
+
+void init_nonbonded_qp_force_kernel_data();
+void calc_nonbonded_qp_forces_host_v2();
+
+void cleanup_nonbonded_qp_force();

src/core/cuda/include/CudaNonbondedQQForce.cuh

@@ -1,8 +1,7 @@
-#ifndef CUDA_NONBONDED_QQ_FORCE_H
-#define CUDA_NONBONDED_QQ_FORCE_H
+#pragma once
 #include "system.h"
 
+void init_nonbonded_qq_force_kernel_data();
 void calc_nonbonded_qq_forces_host();
-void cleanup_nonbonded_qq_force();
 
-#endif
+void cleanup_nonbonded_qq_force();

src/core/cuda/include/CudaNonbondedQWForce.cuh

@@ -0,0 +1,7 @@
+#pragma once
+#include "system.h"
+void init_nonbonded_qw_force_kernel_data();
+
+void calc_nonbonded_qw_forces_host_v2();
+
+void cleanup_nonbonded_qw_force();

src/core/cuda/include/CudaNonbondedWWForce.cuh

@@ -0,0 +1,7 @@
+#pragma once
+#include "system.h"
+void init_nonbonded_ww_force_kernel_data();
+
+void calc_nonbonded_ww_forces_host_v2();
+
+void cleanup_nonbonded_ww_force();

src/core/cuda/include/CudaPolxWaterForce.cuh

@@ -1,10 +1,8 @@
-#ifndef __CUDA_POLX_WATER_FORCE_CUH__
-#define __CUDA_POLX_WATER_FORCE_CUH__
-
+#pragma once
 #include "system.h"
 
+void init_polx_water_force_kernel_data();
+
 void calc_polx_water_forces_host(int iteration);
 
 void cleanup_polx_water_force();
-
-#endif

src/core/cuda/include/CudaPshellForce.cuh

@@ -1,17 +1,7 @@
-#ifndef CUDA_PSHELL_FORCE_CUH
-#define CUDA_PSHELL_FORCE_CUH
+#pragma once
 #include "system.h"
-__global__ void calc_pshell_force_kernel(
-    int n_atoms_solute,
-    bool* shell,
-    bool* excluded,
-    coord_t* coords,
-    coord_t* coords_top,
-    double* ufix_energy,
-    double* ushell_energy,
-    dvel_t* dvelocities
-);
+void init_pshell_force_kernel_data();
+
 void calc_pshell_forces_host();
-void cleanup_pshell_force();
 
-#endif  // CUDA_PSHELL_FORCE_CUH
+void cleanup_pshell_force();

src/core/cuda/include/CudaRadixWaterForce.cuh

@@ -1,22 +1,6 @@
-#ifndef CUDA_RADIX_WATER_FORCE_CU
-#define CUDA_RADIX_WATER_FORCE_CU
+#pragma once
 #include "system.h"
 
-
-__global__ void calc_radix_water_forces_kernel(
-    coord_t* coords, 
-    double shift, 
-    int n_atoms_solute,
-    int n_atoms,
-    topo_t topo,
-    md_t md,
-    double Dwmz,
-    double awmz,
-    dvel_t* dvelocities, 
-    double* energy);
-
+void init_radix_water_force_kernel_data();
 void calc_radix_water_forces_host();
-
 void cleanup_radix_water_force();
-
-#endif

src/core/cuda/include/CudaRestrangForce.cuh

@@ -1,8 +1,7 @@
-#ifndef CUDA_RESTRANG_FORCE_CUH
-#define CUDA_RESTRANG_FORCE_CUH
+#pragma once
 #include "system.h"
 
+void init_restrang_force_kernel_data();
 void calc_restrang_force_host();
-void cleanup_restrang_force();
 
-#endif  // CUDA_RESTRANG_FORCE_CUH
+void cleanup_restrang_force();

src/core/cuda/include/CudaRestrdisForce.cuh

@@ -1,10 +1,7 @@
-#ifndef CUDA_RESTRDIS_FORCE_CUH
-#define CUDA_RESTRDIS_FORCE_CUH
+#pragma once
 #include "system.h"
-__global__ void calc_restrdis_forces_kernel();
 
+void init_restrdis_force_kernel_data();
 void calc_restrdis_forces_host();
 
 void cleanup_restrdis_force();
-
-#endif

src/core/cuda/include/CudaRestrposForce.cuh

@@ -1,8 +1,7 @@
-#ifndef CUDA_RESTRPOS_FORCE_CUH
-#define CUDA_RESTRPOS_FORCE_CUH
+#pragma once
 #include "system.h"
 
+void init_restrpos_force_kernel_data();
 void calc_restrpos_forces_host();
-void cleanup_restrpos_force();
 
-#endif  // CUDA_RESTRPOS_FORCE_CUH
+void cleanup_restrpos_force();

src/core/cuda/include/CudaRestrseqForce.cuh

@@ -1,11 +1,6 @@
-#ifndef CUDA_RESTRSEQ_FORCE_CUH
-#define CUDA_RESTRSEQ_FORCE_CUH
+#pragma once
 #include "system.h"
-__global__ void calc_restrseq_forces_kernel();
 
+void init_restrseq_force_kernel_data();
 void calc_restrseq_forces_host();
-
 void cleanup_restrseq_force();
-
-
-#endif // CUDA_RESTRSEQ_FORCE_CUH

src/core/cuda/include/CudaRestrwallForce.cuh

@@ -1,8 +1,6 @@
-#ifndef CUDA_RESTRWALL_FORCE_CUH
-#define CUDA_RESTRWALL_FORCE_CUH
+#pragma once
 #include "system.h"
 
+void init_restrwall_force_kernel_data();
 void calc_restrwall_forces_host();
 void cleanup_restrwall_force();
-
-#endif