From 7f068a69595e35c83c5c2477ac6f0ba9bc545479 Mon Sep 17 00:00:00 2001
From: Martin Bauernfeind <martin.m.bauernfeind@fau.de>
Date: Thu, 23 Jun 2022 19:32:09 +0200
Subject: [PATCH] :recycle: Fixing refactoring step

---
 src/force.cu | 20 ++++++++++----------
 src/main.c   |  4 ++--
 2 files changed, 12 insertions(+), 12 deletions(-)

diff --git a/src/force.cu b/src/force.cu
index c5aa838..f7ef307 100644
--- a/src/force.cu
+++ b/src/force.cu
@@ -40,7 +40,7 @@ extern "C" {
 
 // cuda kernel
 __global__ void calc_force(
-    Atom a,
+    Atom *a,
     MD_FLOAT cutforcesq, MD_FLOAT sigma6, MD_FLOAT epsilon,
     int Nlocal, int neigh_maxneighs, int *neigh_neighbors, int *neigh_numneigh) {
 
@@ -49,7 +49,7 @@ __global__ void calc_force(
         return;
     }
 
-    Atom *atom = &a;
+    Atom *atom = a;
 
     const int numneighs = neigh_numneigh[i];
 
@@ -139,7 +139,7 @@ int get_num_threads() {
     return num_threads;
 }
 
-void cuda_final_integrate(bool doReneighbour, Parameter *param, Atom *atom) {
+void cuda_final_integrate(bool doReneighbour, Parameter *param, Atom *atom, Atom *c_atom) {
 
     const int Nlocal = atom->Nlocal;
     const int num_threads = get_num_threads();
@@ -157,7 +157,7 @@ void cuda_final_integrate(bool doReneighbour, Parameter *param, Atom *atom) {
     }
 }
 
-void cuda_initial_integrate(bool doReneighbour, Parameter *param, Atom *atom) {
+void cuda_initial_integrate(bool doReneighbour, Parameter *param, Atom *atom, Atom *c_atom) {
 
     const int Nlocal = atom->Nlocal;
     const int num_threads = get_num_threads();
@@ -194,11 +194,11 @@ double computeForce(
 
     const int num_threads = get_num_threads();
 
-    c_atom.Natoms = atom->Natoms;
-    c_atom.Nlocal = atom->Nlocal;
-    c_atom.Nghost = atom->Nghost;
-    c_atom.Nmax = atom->Nmax;
-    c_atom.ntypes = atom->ntypes;
+    c_atom->Natoms = atom->Natoms;
+    c_atom->Nlocal = atom->Nlocal;
+    c_atom->Nghost = atom->Nghost;
+    c_atom->Nmax = atom->Nmax;
+    c_atom->ntypes = atom->ntypes;
 
     /*
     int nDevices;
@@ -219,7 +219,7 @@ double computeForce(
 
     cudaProfilerStart();
 
-    checkCUDAError( "c_atom.x memcpy", cudaMemcpy(c_atom.x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) );
+    checkCUDAError( "c_atom->x memcpy", cudaMemcpy(c_atom->x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) );
 
     if(reneighbourHappenend) {
         checkCUDAError( "c_neighbor->numneigh memcpy", cudaMemcpy(c_neighbor->numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) );
diff --git a/src/main.c b/src/main.c
index 2aeb545..714adc7 100644
--- a/src/main.c
+++ b/src/main.c
@@ -310,7 +310,7 @@ int main(int argc, char** argv)
 
         const bool doReneighbour = (n + 1) % param.every == 0;
 
-        cuda_initial_integrate(doReneighbour, &param, &atom);
+        cuda_initial_integrate(doReneighbour, &param, &atom, &c_atom);
 
         if(doReneighbour) {
             timer[NEIGH] += reneighbour(&param, &atom, &neighbor);
@@ -328,7 +328,7 @@ int main(int argc, char** argv)
 #endif
         }
 
-        cuda_final_integrate(doReneighbour, &param, &atom);
+        cuda_final_integrate(doReneighbour, &param, &atom, &c_atom);
 
         if(!((n + 1) % param.nstat) && (n+1) < param.ntimes) {
             computeThermo(n + 1, &param, &atom);