Adjustments for coherence between ParticleCount updates on CPU and GP…

…U. Use _particle->_uid as identifier in the lists
McStasMcXtrace · Jan 15, 2025 · 5ff5e0c · 5ff5e0c
1 parent 4acdb99
commit 5ff5e0c
Showing 1 changed file with 9 additions and 19 deletions.
diff --git a/mcstas-comps/share/monitor_nd-lib.c b/mcstas-comps/share/monitor_nd-lib.c
@@ -992,14 +992,13 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
   long    While_Buffer=0;
   char    Set_Vars_Coord_Type = DEFS->COORD_NONE;
 
-  #ifdef OPENACC
-  /* For the OPENACC list buffer we need a local copy of the
-     atomically updated Neutron_counter - captured below under list mode */
-  long long ParticleCount=0;
-  #else
-  /* On CPU can make do with the global, declared one... */
-  #define ParticleCount Vars->Neutron_Counter
-  #endif
+  /* For the OPENACC list buffer an atomic capture/update of the
+     updated Neutron_counter - captured below under list mode */
+  long long ParticleCount;
+  #pragma acc atomic capture
+  {
+    ParticleCount=Vars->Neutron_Counter++;
+  }
 
   /* the logic below depends mainly on:
        Flag_List:        1=store 1 buffer, 2=list all, 3=re-use buffer 
@@ -1042,8 +1041,7 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
       Vars->Flag_List = 3;
       Vars->Buffer_Block = Vars->Buffer_Size;
       Vars->Buffer_Counter  = 0;
-      ParticleCount = 0;
-
+      Vars->Neutron_Counter = 0;
     }
     else
     {
@@ -1255,7 +1253,7 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
         else
         if (Set_Vars_Coord_Type == DEFS->COORD_LAMBDA) { XY = sqrt(_particle->vx*_particle->vx+_particle->vy*_particle->vy+_particle->vz*_particle->vz);  XY *= V2K; if (XY != 0) XY = 2*PI/XY; }
         else
-        if (Set_Vars_Coord_Type == DEFS->COORD_NCOUNT) XY = ParticleCount;
+	  if (Set_Vars_Coord_Type == DEFS->COORD_NCOUNT) XY = _particle->_uid;
         else
         if (Set_Vars_Coord_Type == DEFS->COORD_ANGLE)
         {  XY = sqrt(_particle->vx*_particle->vx+_particle->vy*_particle->vy);
@@ -1382,14 +1380,6 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
     { /* now store Coord into Buffer (no index needed) if necessary (list or auto limits) */
       if ((Vars->Buffer_Counter < Vars->Buffer_Block) && ((Vars->Flag_List) || (Vars->Flag_Auto_Limits == 1)))
       {
-        #ifdef OPENACC
-        #pragma acc atomic capture
-        {
-	  ParticleCount=Vars->Neutron_Counter++ ;
-        } 
-        #else
-        ParticleCount++;
-	#endif
         for (i = 0; i <= Vars->Coord_Number; i++)
         {
 	  // This is is where the list is appended. How to make this "atomic"?