From 0e92ae26deddfad043bddaf883ebf738a5949de1 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 15:02:12 +0800
Subject: [PATCH 01/23] calculate stopping_power_loss

---
 src/main_gpumd/electron_stop.cu  | 4 +++-
 src/main_gpumd/electron_stop.cuh | 1 +
 src/main_gpumd/run.cu            | 2 +-
 3 files changed, 5 insertions(+), 2 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 43c6f7b7d..a24853039 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -66,6 +66,8 @@ static void __global__ find_stopping_force(
     g_force[0 * num_atoms + i] = vx * factor;
     g_force[1 * num_atoms + i] = vy * factor;
     g_force[2 * num_atoms + i] = vz * factor;
+
+    stopping_power_loss += stopping_power * sqrt(v2) * time_step;
   }
 }
 
@@ -119,7 +121,7 @@ apply_electron_stopping(const int num_atoms, const double* g_stopping_force, dou
   }
 }
 
-void Electron_Stop::compute(Atom& atom)
+void Electron_Stop::compute(double time_step, Atom& atom)
 {
   if (!do_electron_stop) {
     return;
diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index 653f92bef..ef3edfd65 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -30,6 +30,7 @@ public:
 
 private:
   int num_points = 0;
+  double stopping_power_loss = 0.0;
   double energy_min;
   double energy_max;
   double energy_interval;
diff --git a/src/main_gpumd/run.cu b/src/main_gpumd/run.cu
index 20c63acac..982971334 100644
--- a/src/main_gpumd/run.cu
+++ b/src/main_gpumd/run.cu
@@ -241,7 +241,7 @@ void Run::perform_a_run()
     }
 #endif
 
-    electron_stop.compute(atom);
+    electron_stop.compute(time_step, atom);
 
     integrate.compute2(time_step, double(step) / number_of_steps, group, box, atom, thermo);
 

From c308ffd7167df7e95cae17786db6190361ec10e5 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 15:52:28 +0800
Subject: [PATCH 02/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 1 -
 1 file changed, 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index ef3edfd65..653f92bef 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -30,7 +30,6 @@ public:
 
 private:
   int num_points = 0;
-  double stopping_power_loss = 0.0;
   double energy_min;
   double energy_max;
   double energy_interval;

From 4cf5b21eb883cce602f5529a6310404cf8b6f2d7 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 15:58:06 +0800
Subject: [PATCH 03/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index a24853039..9f201e2e4 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -27,6 +27,7 @@ Apply electron stopping.
 static void __global__ find_stopping_force(
   const int num_atoms,
   const int num_points,
+  const double time_step,
   const double energy_min,
   const double energy_max,
   const double energy_interval_inverse,
@@ -37,6 +38,7 @@ static void __global__ find_stopping_force(
   double* g_force)
 {
   const int i = blockIdx.x * blockDim.x + threadIdx.x;
+  double stopping_power_loss = 0.0;
   if (i < num_atoms) {
     int type = g_type[i];
     double mass = g_mass[i];

From c09973bc977e72a975e244f5d3fbda9283418d18 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 16:03:53 +0800
Subject: [PATCH 04/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index 653f92bef..ef3edfd65 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -30,6 +30,7 @@ public:
 
 private:
   int num_points = 0;
+  double stopping_power_loss = 0.0;
   double energy_min;
   double energy_max;
   double energy_interval;

From 02dd1f58a3129afbb65fc67f88a16a06e529b595 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 16:03:59 +0800
Subject: [PATCH 05/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 10 ++++++++--
 1 file changed, 8 insertions(+), 2 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 9f201e2e4..b73818422 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -38,7 +38,6 @@ static void __global__ find_stopping_force(
   double* g_force)
 {
   const int i = blockIdx.x * blockDim.x + threadIdx.x;
-  double stopping_power_loss = 0.0;
   if (i < num_atoms) {
     int type = g_type[i];
     double mass = g_mass[i];
@@ -210,4 +209,11 @@ void Electron_Stop::parse(
   do_electron_stop = true;
 }
 
-void Electron_Stop::finalize() { do_electron_stop = false; }
+void Electron_Stop::finalize() 
+{ 
+  if (do_electron_stop) { 
+    printf("    total electron stopping power loss = %g eV.\n", stopping_power_loss);
+  }
+  do_electron_stop = false; 
+  stopping_power_loss = 0.0;
+  }

From d4b16cb4933e4750ae3dbd489cdb68afe221c593 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 16:05:43 +0800
Subject: [PATCH 06/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index b73818422..5f244fed7 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -131,6 +131,7 @@ void Electron_Stop::compute(double time_step, Atom& atom)
   find_stopping_force<<<(atom.number_of_atoms - 1) / 64 + 1, 64>>>(
     atom.number_of_atoms,
     num_points,
+    time_step,
     energy_min,
     energy_max,
     1.0 / energy_interval,

From 785e286c3e10974f031b5384eac81882abd9e36d Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 16:20:02 +0800
Subject: [PATCH 07/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index ef3edfd65..6f76323f3 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -25,7 +25,7 @@ class Electron_Stop
 public:
   bool do_electron_stop = false;
   void parse(const char** param, int num_param, const int num_atoms, const int num_types);
-  void compute(Atom& atom);
+  void compute(double time_step, Atom& atom);
   void finalize();
 
 private:

From 4cc344a1316695e80b931241893962aca552f05a Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 16:26:53 +0800
Subject: [PATCH 08/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index 6f76323f3..d3439f1a1 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -24,13 +24,13 @@ class Electron_Stop
 {
 public:
   bool do_electron_stop = false;
+  double stopping_power_loss = 0.0;
   void parse(const char** param, int num_param, const int num_atoms, const int num_types);
   void compute(double time_step, Atom& atom);
   void finalize();
 
 private:
   int num_points = 0;
-  double stopping_power_loss = 0.0;
   double energy_min;
   double energy_max;
   double energy_interval;

From fe4de1e77895443ddfd31b2f85ad86d72ad8dd4c Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 16:45:39 +0800
Subject: [PATCH 09/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 5f244fed7..2799af0f3 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -35,7 +35,8 @@ static void __global__ find_stopping_force(
   const int* g_type,
   const double* g_mass,
   const double* g_velocity,
-  double* g_force)
+  double* g_force,
+  double* stopping_power_loss)
 {
   const int i = blockIdx.x * blockDim.x + threadIdx.x;
   if (i < num_atoms) {
@@ -139,7 +140,9 @@ void Electron_Stop::compute(double time_step, Atom& atom)
     atom.type.data(),
     atom.mass.data(),
     atom.velocity_per_atom.data(),
-    stopping_force.data());
+    stopping_force.data(),
+    stopping_power_loss);
+
   CUDA_CHECK_KERNEL
 
   find_force_average<<<3, 1024>>>(atom.number_of_atoms, stopping_force.data());
@@ -217,4 +220,4 @@ void Electron_Stop::finalize()
   }
   do_electron_stop = false; 
   stopping_power_loss = 0.0;
-  }
+}

From 0240b1a89c6235af0bee7bd8a697d85001be6715 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:02:31 +0800
Subject: [PATCH 10/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 2799af0f3..34d1e17b9 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -36,7 +36,7 @@ static void __global__ find_stopping_force(
   const double* g_mass,
   const double* g_velocity,
   double* g_force,
-  double* stopping_power_loss)
+  double* g_power_loss)
 {
   const int i = blockIdx.x * blockDim.x + threadIdx.x;
   if (i < num_atoms) {
@@ -69,7 +69,7 @@ static void __global__ find_stopping_force(
     g_force[1 * num_atoms + i] = vy * factor;
     g_force[2 * num_atoms + i] = vz * factor;
 
-    stopping_power_loss += stopping_power * sqrt(v2) * time_step;
+    g_power_loss[i] = stopping_power * sqrt(v2) * time_step;
   }
 }
 
@@ -141,7 +141,7 @@ void Electron_Stop::compute(double time_step, Atom& atom)
     atom.mass.data(),
     atom.velocity_per_atom.data(),
     stopping_force.data(),
-    stopping_power_loss);
+    stopping_power_loss.data();
 
   CUDA_CHECK_KERNEL
 
@@ -210,6 +210,7 @@ void Electron_Stop::parse(
   stopping_power_gpu.resize(num_points * num_types);
   stopping_power_gpu.copy_from_host(stopping_power_cpu.data());
   stopping_force.resize(num_atoms * 3);
+  stopping_power_loss.resize(num_atoms);
   do_electron_stop = true;
 }
 

From 6cee69790c78b1edbff6c0360595ecc40834126e Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:02:35 +0800
Subject: [PATCH 11/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index d3439f1a1..40ad1a264 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -37,4 +37,5 @@ private:
   std::vector<double> stopping_power_cpu;
   GPU_Vector<double> stopping_power_gpu;
   GPU_Vector<double> stopping_force;
+  GPU_Vector<double> stopping_power_lose;
 };

From a7ab5c6b2b4e2359c47c32a691b91e9a61f933ef Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:27:22 +0800
Subject: [PATCH 12/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index 40ad1a264..0db5f9100 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -37,5 +37,5 @@ private:
   std::vector<double> stopping_power_cpu;
   GPU_Vector<double> stopping_power_gpu;
   GPU_Vector<double> stopping_force;
-  GPU_Vector<double> stopping_power_lose;
+  GPU_Vector<double> stopping_power_loss;
 };

From 768efb006a5839f6707c76d93ac3e9f4ecda0f10 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:27:28 +0800
Subject: [PATCH 13/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 34d1e17b9..7084c47a5 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -141,7 +141,7 @@ void Electron_Stop::compute(double time_step, Atom& atom)
     atom.mass.data(),
     atom.velocity_per_atom.data(),
     stopping_force.data(),
-    stopping_power_loss.data();
+    stopping_power_loss.data());
 
   CUDA_CHECK_KERNEL
 

From c496a74d7e6553525d382d65429acded002754de Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:43:02 +0800
Subject: [PATCH 14/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 7084c47a5..2699ff48d 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -69,7 +69,8 @@ static void __global__ find_stopping_force(
     g_force[1 * num_atoms + i] = vy * factor;
     g_force[2 * num_atoms + i] = vz * factor;
 
-    g_power_loss[i] = stopping_power * sqrt(v2) * time_step;
+    power_loss = stopping_power * sqrt(v2) * time_step;
+    atomicAdd(g_power_loss, power_loss);
   }
 }
 
@@ -210,7 +211,7 @@ void Electron_Stop::parse(
   stopping_power_gpu.resize(num_points * num_types);
   stopping_power_gpu.copy_from_host(stopping_power_cpu.data());
   stopping_force.resize(num_atoms * 3);
-  stopping_power_loss.resize(num_atoms);
+  
   do_electron_stop = true;
 }
 

From 64443e9a37d6d7ed745214e86ce54df734b17c21 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:43:06 +0800
Subject: [PATCH 15/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 1 -
 1 file changed, 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index 0db5f9100..d3439f1a1 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -37,5 +37,4 @@ private:
   std::vector<double> stopping_power_cpu;
   GPU_Vector<double> stopping_power_gpu;
   GPU_Vector<double> stopping_force;
-  GPU_Vector<double> stopping_power_loss;
 };

From 31e788ff9e36096cf26f04ecf784126d5cd72e29 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 17:45:37 +0800
Subject: [PATCH 16/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 2699ff48d..60f15e349 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -141,8 +141,7 @@ void Electron_Stop::compute(double time_step, Atom& atom)
     atom.type.data(),
     atom.mass.data(),
     atom.velocity_per_atom.data(),
-    stopping_force.data(),
-    stopping_power_loss.data());
+    stopping_force.data());
 
   CUDA_CHECK_KERNEL
 
@@ -211,7 +210,6 @@ void Electron_Stop::parse(
   stopping_power_gpu.resize(num_points * num_types);
   stopping_power_gpu.copy_from_host(stopping_power_cpu.data());
   stopping_force.resize(num_atoms * 3);
-  
   do_electron_stop = true;
 }
 

From 8ad8bf7ee9bdd80d52ac5b29c4cb59f22a25ba68 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 19:59:30 +0800
Subject: [PATCH 17/23] Update electron_stop.cuh

---
 src/main_gpumd/electron_stop.cuh | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/main_gpumd/electron_stop.cuh b/src/main_gpumd/electron_stop.cuh
index d3439f1a1..912181ce3 100644
--- a/src/main_gpumd/electron_stop.cuh
+++ b/src/main_gpumd/electron_stop.cuh
@@ -37,4 +37,5 @@ private:
   std::vector<double> stopping_power_cpu;
   GPU_Vector<double> stopping_power_gpu;
   GPU_Vector<double> stopping_force;
+  GPU_Vector<double> stopping_loss;
 };

From 8a6e5d30d8742da162de493ca991506f0d4ab999 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 19:59:44 +0800
Subject: [PATCH 18/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 60f15e349..2295863c8 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -69,8 +69,7 @@ static void __global__ find_stopping_force(
     g_force[1 * num_atoms + i] = vy * factor;
     g_force[2 * num_atoms + i] = vz * factor;
 
-    power_loss = stopping_power * sqrt(v2) * time_step;
-    atomicAdd(g_power_loss, power_loss);
+    g_power_loss[i] = stopping_power * sqrt(v2) * time_step; 
   }
 }
 
@@ -210,6 +209,7 @@ void Electron_Stop::parse(
   stopping_power_gpu.resize(num_points * num_types);
   stopping_power_gpu.copy_from_host(stopping_power_cpu.data());
   stopping_force.resize(num_atoms * 3);
+  stopping_loss.resize(num_atoms);
   do_electron_stop = true;
 }
 

From 31f0c82dfbd8dbb9106f5bf71200132d290cf207 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 20:18:42 +0800
Subject: [PATCH 19/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 50 ++++++++++++++++++++++++++++++++-
 1 file changed, 49 insertions(+), 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 2295863c8..4b9134475 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -123,6 +123,43 @@ apply_electron_stopping(const int num_atoms, const double* g_stopping_force, dou
   }
 }
 
+__device__ float device_power_loss;
+
+static __global__ void find_power_loss(int num_atoms, double* g_power_loss)
+{
+  //<<<1, 1024>>>
+  int tid = threadIdx.x;
+  int bid = blockIdx.x;
+    int num_blocks = gridDim.x;
+    int block_size = blockDim.x;
+
+    int number_of_batches = (num_atoms + block_size - 1) / block_size;
+    __shared__ double s_f[1024];
+    double f = 0.0;
+
+    for (int batch = 0; batch < number_of_batches; ++batch) {
+        int idx = tid + batch * block_size;
+        if (idx < num_atoms) {
+            f += g_power_loss[idx];
+        }
+    }
+
+    s_f[tid] = f;
+    __syncthreads();
+
+    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1) {
+        if (tid < offset) {
+            s_f[tid] += s_f[tid + offset];
+        }
+        __syncthreads(); 
+    }
+
+    if (tid == 0) {
+        device_power_loss = s_f[0];
+    }
+
+}
+
 void Electron_Stop::compute(double time_step, Atom& atom)
 {
   if (!do_electron_stop) {
@@ -140,7 +177,8 @@ void Electron_Stop::compute(double time_step, Atom& atom)
     atom.type.data(),
     atom.mass.data(),
     atom.velocity_per_atom.data(),
-    stopping_force.data());
+    stopping_force.data(),
+    stopping_loss.data());
 
   CUDA_CHECK_KERNEL
 
@@ -150,6 +188,16 @@ void Electron_Stop::compute(double time_step, Atom& atom)
   apply_electron_stopping<<<(atom.number_of_atoms - 1) / 64 + 1, 64>>>(
     atom.number_of_atoms, stopping_force.data(), atom.force_per_atom.data());
   CUDA_CHECK_KERNEL
+
+  find_power_loss<<<1, 1024>>>(atom.number_of_atoms, stopping_loss.data());
+  CUDA_CHECK_KERNEL
+
+  cudaDeviceSynchronize();
+  double power_loss_host;
+  cudaMemcpyFromSymbol(&power_loss_host, device_power_loss, sizeof(double), 0, cudaMemcpyDeviceToHost);
+  stopping_power_loss += power_loss_host;
+  cudaMemset(&device_power_loss, 0, sizeof(double));
+
 }
 
 void Electron_Stop::parse(

From 08db84d0f7e05d41242c1c6aa92d447b54e0459d Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 20:27:12 +0800
Subject: [PATCH 20/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 48 ++++++++++++++++-----------------
 1 file changed, 23 insertions(+), 25 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 4b9134475..8a49551e6 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -129,34 +129,32 @@ static __global__ void find_power_loss(int num_atoms, double* g_power_loss)
 {
   //<<<1, 1024>>>
   int tid = threadIdx.x;
-  int bid = blockIdx.x;
-    int num_blocks = gridDim.x;
-    int block_size = blockDim.x;
-
-    int number_of_batches = (num_atoms + block_size - 1) / block_size;
-    __shared__ double s_f[1024];
-    double f = 0.0;
-
-    for (int batch = 0; batch < number_of_batches; ++batch) {
-        int idx = tid + batch * block_size;
-        if (idx < num_atoms) {
-            f += g_power_loss[idx];
-        }
-    }
+  int block_size = blockDim.x;
 
-    s_f[tid] = f;
-    __syncthreads();
+  int number_of_batches = (num_atoms + block_size - 1) / block_size;
+  __shared__ double s_f[1024];
+  double f = 0.0;
 
-    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1) {
-        if (tid < offset) {
-            s_f[tid] += s_f[tid + offset];
-        }
-        __syncthreads(); 
-    }
+  for (int batch = 0; batch < number_of_batches; ++batch) {
+      int idx = tid + batch * block_size;
+      if (idx < num_atoms) {
+          f += g_power_loss[idx];
+      }
+  }
 
-    if (tid == 0) {
-        device_power_loss = s_f[0];
-    }
+  s_f[tid] = f;
+  __syncthreads();
+
+  for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1) {
+      if (tid < offset) {
+          s_f[tid] += s_f[tid + offset];
+      }
+      __syncthreads(); 
+  }
+
+  if (tid == 0) {
+      device_power_loss = s_f[0];
+  }
 
 }
 

From e22f3f74556be9acc5f76b121dce58533f23846a Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Tue, 28 May 2024 21:55:32 +0800
Subject: [PATCH 21/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 9 +++------
 1 file changed, 3 insertions(+), 6 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index 8a49551e6..e57fd2b5a 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -190,12 +190,9 @@ void Electron_Stop::compute(double time_step, Atom& atom)
   find_power_loss<<<1, 1024>>>(atom.number_of_atoms, stopping_loss.data());
   CUDA_CHECK_KERNEL
 
-  cudaDeviceSynchronize();
-  double power_loss_host;
-  cudaMemcpyFromSymbol(&power_loss_host, device_power_loss, sizeof(double), 0, cudaMemcpyDeviceToHost);
-  stopping_power_loss += power_loss_host;
-  cudaMemset(&device_power_loss, 0, sizeof(double));
-
+  float power_loss_host;  
+  cudaMemcpyFromSymbol(&power_loss_host, device_power_loss, sizeof(float), 0, cudaMemcpyDeviceToHost);
+  stopping_power_loss += static_cast<double>(power_loss_host);
 }
 
 void Electron_Stop::parse(

From ba9cdedb2b2f97d20104e6662f735376bab42723 Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Wed, 29 May 2024 14:08:59 +0800
Subject: [PATCH 22/23] change float to double

---
 src/main_gpumd/electron_stop.cu | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index e57fd2b5a..febd917ba 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -123,7 +123,7 @@ apply_electron_stopping(const int num_atoms, const double* g_stopping_force, dou
   }
 }
 
-__device__ float device_power_loss;
+__device__ double device_power_loss;
 
 static __global__ void find_power_loss(int num_atoms, double* g_power_loss)
 {
@@ -190,9 +190,9 @@ void Electron_Stop::compute(double time_step, Atom& atom)
   find_power_loss<<<1, 1024>>>(atom.number_of_atoms, stopping_loss.data());
   CUDA_CHECK_KERNEL
 
-  float power_loss_host;  
-  cudaMemcpyFromSymbol(&power_loss_host, device_power_loss, sizeof(float), 0, cudaMemcpyDeviceToHost);
-  stopping_power_loss += static_cast<double>(power_loss_host);
+  double power_loss_host;  
+  CHECK(cudaMemcpyFromSymbol(&power_loss_host, device_power_loss, sizeof(double), 0, cudaMemcpyDeviceToHost));
+  stopping_power_loss += power_loss_host;
 }
 
 void Electron_Stop::parse(

From dad95d4ebe272d911ef4a917aee471cf87e58a7b Mon Sep 17 00:00:00 2001
From: Jonsnow-willow <55383669+Jonsnow-willow@users.noreply.github.com>
Date: Wed, 5 Jun 2024 15:35:31 +0800
Subject: [PATCH 23/23] Update electron_stop.cu

---
 src/main_gpumd/electron_stop.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main_gpumd/electron_stop.cu b/src/main_gpumd/electron_stop.cu
index febd917ba..3f07d0397 100644
--- a/src/main_gpumd/electron_stop.cu
+++ b/src/main_gpumd/electron_stop.cu
@@ -259,7 +259,7 @@ void Electron_Stop::parse(
 void Electron_Stop::finalize() 
 { 
   if (do_electron_stop) { 
-    printf("    total electron stopping power loss = %g eV.\n", stopping_power_loss);
+    printf("Total electron stopping power loss = %g eV.\n", stopping_power_loss);
   }
   do_electron_stop = false; 
   stopping_power_loss = 0.0;