Core: fixes for CUDA usage of Backend_par.

spirit-code · Apr 20, 2020 · 3228c39 · 3228c39
1 parent 8ed4aef
commit 3228c39
Show file tree

Hide file tree

Showing 2 changed files with 15 additions and 13 deletions.
diff --git a/core/include/engine/Hamiltonian_Heisenberg.hpp b/core/include/engine/Hamiltonian_Heisenberg.hpp
@@ -152,11 +152,12 @@ namespace Engine
         // Calculate the Quadruplet energy
         void E_Quadruplet(const vectorfield & spins, scalarfield & Energy);
 
-        private:
-        int idx_zeeman, idx_anisotropy, idx_exchange, idx_dmi, idx_ddi, idx_quadruplet;
         void Gradient_DDI_Cutoff(const vectorfield& spins, vectorfield & gradient);
         void Gradient_DDI_Direct(const vectorfield& spins, vectorfield & gradient);
         void Gradient_DDI_FFT(const vectorfield& spins, vectorfield & gradient);
+
+    private:
+        int idx_zeeman, idx_anisotropy, idx_exchange, idx_dmi, idx_ddi, idx_quadruplet;
         void E_DDI_Direct(const vectorfield& spins, scalarfield & Energy);
         void E_DDI_Cutoff(const vectorfield& spins, scalarfield & Energy);
         void E_DDI_FFT(const vectorfield& spins, scalarfield & Energy);

diff --git a/core/src/engine/Hamiltonian_Heisenberg.cu b/core/src/engine/Hamiltonian_Heisenberg.cu
@@ -228,7 +228,7 @@ namespace Engine
             else
             {
                 for (auto& pair : contributions)
-                    Backend::par::assign(pair.second, [] SPIRIT_LAMBDA (int idx) {return 0;});
+                    Backend::par::assign(pair.second, [] SPIRIT_LAMBDA () {return 0;});
             }
         }
 
@@ -371,12 +371,12 @@ namespace Engine
     void Hamiltonian_Heisenberg::E_DDI_Direct(const vectorfield & spins, scalarfield & Energy)
     {
         vectorfield gradients_temp;
-        gradients_temp.resize(geometry->nos);
-        Backend::par::assign(gradients_temp, [] SPIRIT_LAMBDA (int idx) -> Vector3 {return {0, 0, 0};});
+        gradients_temp.resize(spins.size());
+        Backend::par::assign(gradients_temp, [] SPIRIT_LAMBDA () -> Vector3 {return {0, 0, 0};});
         this->Gradient_DDI_Direct(spins, gradients_temp);
 
         #pragma omp parallel for
-        for (int ispin = 0; ispin < geometry->nos; ispin++)
+        for (int ispin = 0; ispin < spins.size(); ispin++)
         {
             Energy[ispin] += 0.5 * spins[ispin].dot(gradients_temp[ispin]);
         }
@@ -424,17 +424,18 @@ namespace Engine
     }
     void Hamiltonian_Heisenberg::E_DDI_FFT(const vectorfield & spins, scalarfield & Energy)
     {
+        auto N = spins.size();
         //todo maybe the gradient should be cached somehow, it is quite inefficient to calculate it
         //again just for the energy
-        vectorfield gradients_temp(geometry->nos);
-        Backend::par::assign(gradients_temp, [] SPIRIT_LAMBDA (int idx) -> Vector3 {return {0, 0, 0};});
-        this->Gradient_DDI(spins, gradients_temp);
-        CU_E_DDI_FFT<<<(geometry->nos + 1023)/1024, 1024>>>(Energy.data(), spins.data(), gradients_temp.data(), geometry->nos, geometry->n_cell_atoms, geometry->mu_s.data());
+        vectorfield gradients_temp(N);
+        Backend::par::assign(gradients_temp, [] SPIRIT_LAMBDA () -> Vector3 {return {0, 0, 0};});
+        this->Gradient_DDI_FFT(spins, gradients_temp);
+        CU_E_DDI_FFT<<<(N + 1023)/1024, 1024>>>(Energy.data(), spins.data(), gradients_temp.data(), N, geometry->n_cell_atoms, geometry->mu_s.data());
 
         // === DEBUG: begin gradient comparison ===
             // vectorfield gradients_temp_dir;
             // gradients_temp_dir.resize(this->geometry->nos);
-            // Backend::par::assign(gradients_temp_dir, [] SPIRIT_LAMBDA (int idx) -> Vector3 {return {0, 0, 0};});
+            // Backend::par::assign(gradients_temp_dir, [] SPIRIT_LAMBDA () -> Vector3 {return {0, 0, 0};});
             // Gradient_DDI_Direct(spins, gradients_temp_dir);
 
             // //get deviation
@@ -626,7 +627,7 @@ namespace Engine
     void Hamiltonian_Heisenberg::Gradient(const vectorfield & spins, vectorfield & gradient)
     {
         // Set to zero
-        Backend::par::assign(gradient, [] SPIRIT_LAMBDA (int idx) -> Vector3 {return {0, 0, 0};});
+        Backend::par::assign(gradient, [] SPIRIT_LAMBDA () -> Vector3 {return {0, 0, 0};});
 
         // External field
         if(idx_zeeman >= 0)
@@ -656,7 +657,7 @@ namespace Engine
     void Hamiltonian_Heisenberg::Gradient_and_Energy(const vectorfield & spins, vectorfield & gradient, scalar & energy)
     {
         // Set to zero
-        Backend::par::assign(gradient, [] SPIRIT_LAMBDA (int idx) -> Vector3 {return {0, 0, 0};});
+        Backend::par::assign(gradient, [] SPIRIT_LAMBDA () -> Vector3 {return {0, 0, 0};});
         energy = 0;
 
         auto N = spins.size();