Skip to content

Commit 4662602

Browse files
willendwillend
authored
Merge pull request #1796 from McStasMcXtrace/fix-issue-1794-monitor_nd-atomic-listmode
Fix issue 1794 monitor_nd atomic listmode
2 parents dedcb05 + e02f9e0 commit 4662602

File tree

2 files changed

+48
-17
lines changed

2 files changed

+48
-17
lines changed

‎mcstas-comps/share/monitor_nd-lib.c‎

Lines changed: 24 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -163,7 +163,7 @@ void Monitor_nD_Init(MonitornD_Defines_type *DEFS,
163163
Vars->Coord_NumberNoPixel=0; /* same but without counting PixelID */
164164

165165
Vars->Buffer_Block = MONND_BUFSIZ; /* Buffer size for list or auto limits */
166-
Vars->Neutron_Counter = -1; /* event counter, simulation total counts is mcget_ncount() */
166+
Vars->Neutron_Counter = 0; /* event counter, simulation total counts is mcget_ncount() */
167167
Vars->Buffer_Counter = 0; /* index in Buffer size (for realloc) */
168168
Vars->Buffer_Size = 0;
169169
Vars->He3_pressure = 0;
@@ -992,6 +992,15 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
992992
long While_Buffer=0;
993993
char Set_Vars_Coord_Type = DEFS->COORD_NONE;
994994

995+
#ifdef OPENACC
996+
/* For the OPENACC list buffer we need a local copy of the
997+
atomically updated Neutron_counter - captured below under list mode */
998+
long long ParticleCount=0;
999+
#else
1000+
/* On CPU can make do with the global, declared one... */
1001+
#define ParticleCount Vars->Neutron_Counter
1002+
#endif
1003+
9951004
/* the logic below depends mainly on:
9961005
Flag_List: 1=store 1 buffer, 2=list all, 3=re-use buffer
9971006
Flag_Auto_Limits: 0 (no auto limits/list), 1 (store events into Buffer), 2 (re-emit store events)
@@ -1033,15 +1042,15 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
10331042
Vars->Flag_List = 3;
10341043
Vars->Buffer_Block = Vars->Buffer_Size;
10351044
Vars->Buffer_Counter = 0;
1036-
Vars->Neutron_Counter = 0;
1045+
ParticleCount = 0;
10371046

10381047
}
10391048
else
10401049
{
1041-
Vars->Mon2D_Buffer = (double *)realloc(Vars->Mon2D_Buffer, (Vars->Coord_Number+1)*(Vars->Neutron_Counter+Vars->Buffer_Block)*sizeof(double));
1050+
Vars->Mon2D_Buffer = (double *)realloc(Vars->Mon2D_Buffer, (Vars->Coord_Number+1)*(ParticleCount+Vars->Buffer_Block)*sizeof(double));
10421051
if (Vars->Mon2D_Buffer == NULL)
1043-
{ printf("Monitor_nD: %s cannot reallocate Vars->Mon2D_Buffer[%li] (%li). Skipping.\n", Vars->compcurname, i, (long int)(Vars->Neutron_Counter+Vars->Buffer_Block)*sizeof(double)); Vars->Flag_List = 1; }
1044-
else { Vars->Buffer_Counter = 0; Vars->Buffer_Size = Vars->Neutron_Counter+Vars->Buffer_Block; }
1052+
{ printf("Monitor_nD: %s cannot reallocate Vars->Mon2D_Buffer[%li] (%li). Skipping.\n", Vars->compcurname, i, (long int)(ParticleCount+Vars->Buffer_Block)*sizeof(double)); Vars->Flag_List = 1; }
1053+
else { Vars->Buffer_Counter = 0; Vars->Buffer_Size = ParticleCount+Vars->Buffer_Block; }
10451054
}
10461055
} /* end if Buffer realloc */
10471056
#endif
@@ -1246,7 +1255,7 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
12461255
else
12471256
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; }
12481257
else
1249-
if (Set_Vars_Coord_Type == DEFS->COORD_NCOUNT) XY = Vars->Neutron_Counter;
1258+
if (Set_Vars_Coord_Type == DEFS->COORD_NCOUNT) XY = ParticleCount;
12501259
else
12511260
if (Set_Vars_Coord_Type == DEFS->COORD_ANGLE)
12521261
{ XY = sqrt(_particle->vx*_particle->vx+_particle->vy*_particle->vy);
@@ -1373,13 +1382,19 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
13731382
{ /* now store Coord into Buffer (no index needed) if necessary (list or auto limits) */
13741383
if ((Vars->Buffer_Counter < Vars->Buffer_Block) && ((Vars->Flag_List) || (Vars->Flag_Auto_Limits == 1)))
13751384
{
1376-
#pragma acc atomic
1377-
Vars->Neutron_Counter = Vars->Neutron_Counter + 1;
1385+
#ifdef OPENACC
1386+
#pragma acc atomic capture
1387+
{
1388+
ParticleCount=Vars->Neutron_Counter++ ;
1389+
}
1390+
#else
1391+
ParticleCount++;
1392+
#endif
13781393
for (i = 0; i <= Vars->Coord_Number; i++)
13791394
{
13801395
// This is is where the list is appended. How to make this "atomic"?
13811396
#pragma acc atomic write
1382-
Vars->Mon2D_Buffer[i + Vars->Neutron_Counter*(Vars->Coord_Number+1)] = Coord[i];
1397+
Vars->Mon2D_Buffer[i + ParticleCount*(Vars->Coord_Number+1)] = Coord[i];
13831398
}
13841399
#pragma acc atomic
13851400
Vars->Buffer_Counter = Vars->Buffer_Counter + 1;

‎mcxtrace-comps/share/monitor_nd-lib.c‎

Lines changed: 24 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -759,6 +759,15 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
759759
long While_Buffer=0;
760760
char Set_Vars_Coord_Type = DEFS->COORD_NONE;
761761

762+
#ifdef OPENACC
763+
/* For the OPENACC list buffer we need a local copy of the
764+
atomically updated Neutron_counter - captured below under list mode */
765+
long long ParticleCount=0;
766+
#else
767+
/* On CPU can make do with the global, declared one... */
768+
#define ParticleCount Vars->Photon_Counter
769+
#endif
770+
762771
/* the logic below depends mainly on:
763772
Flag_List: 1=store 1 buffer, 2=list all, 3=re-use buffer
764773
Flag_Auto_Limits: 0 (no auto limits/list), 1 (store events into Buffer), 2 (re-emit store events)
@@ -800,14 +809,15 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
800809
Vars->Flag_List = 3;
801810
Vars->Buffer_Block = Vars->Buffer_Size;
802811
Vars->Buffer_Counter = 0;
803-
Vars->Photon_Counter = 0;
812+
ParticleCount = 0;
813+
804814
}
805815
else
806816
{
807-
Vars->Mon2D_Buffer = (double *)realloc(Vars->Mon2D_Buffer, (Vars->Coord_Number+1)*(Vars->Photon_Counter+Vars->Buffer_Block)*sizeof(double));
817+
Vars->Mon2D_Buffer = (double *)realloc(Vars->Mon2D_Buffer, (Vars->Coord_Number+1)*(ParticleCount+Vars->Buffer_Block)*sizeof(double));
808818
if (Vars->Mon2D_Buffer == NULL)
809-
{ printf("Monitor_nD: %s cannot reallocate Vars->Mon2D_Buffer[%li] (%li). Skipping.\n", Vars->compcurname, i, (Vars->Photon_Counter+Vars->Buffer_Block)*sizeof(double)); Vars->Flag_List = 1; }
810-
else { Vars->Buffer_Counter = 0; Vars->Buffer_Size = Vars->Photon_Counter+Vars->Buffer_Block; }
819+
{ printf("Monitor_nD: %s cannot reallocate Vars->Mon2D_Buffer[%li] (%li). Skipping.\n", Vars->compcurname, i, (long int)(ParticleCount+Vars->Buffer_Block)*sizeof(double)); Vars->Flag_List = 1; }
820+
else { Vars->Buffer_Counter = 0; Vars->Buffer_Size = ParticleCount+Vars->Buffer_Block; }
811821
}
812822
} /* end if Buffer realloc */
813823
#endif
@@ -976,7 +986,7 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
976986
else
977987
if (Set_Vars_Coord_Type == DEFS->COORD_LAMBDA) { if (k!=0) XY = 2*M_PI/k; }
978988
else
979-
if (Set_Vars_Coord_Type == DEFS->COORD_NCOUNT) XY = Vars->Photon_Counter;
989+
if (Set_Vars_Coord_Type == DEFS->COORD_NCOUNT) XY = ParticleCount;
980990
else
981991
if (Set_Vars_Coord_Type == DEFS->COORD_ANGLE)
982992
{ XY = sqrt(_particle->kx*_particle->kx+_particle->ky*_particle->ky);
@@ -1103,19 +1113,25 @@ int Monitor_nD_Trace(MonitornD_Defines_type *DEFS, MonitornD_Variables_type *Var
11031113
{ /* now store Coord into Buffer (no index needed) if necessary (list or auto limits) */
11041114
if ((Vars->Buffer_Counter < Vars->Buffer_Block) && ((Vars->Flag_List) || (Vars->Flag_Auto_Limits == 1)))
11051115
{
1116+
#ifdef OPENACC
1117+
#pragma acc atomic capture
1118+
{
1119+
ParticleCount=Vars->Photon_Counter++ ;
1120+
}
1121+
#else
1122+
ParticleCount++;
1123+
#endif
11061124
for (i = 0; i <= Vars->Coord_Number; i++)
11071125
{
11081126
// This is is where the list is appended. How to make this "atomic"?
11091127
#pragma acc atomic write
1110-
Vars->Mon2D_Buffer[i + Vars->Photon_Counter*(Vars->Coord_Number+1)] = Coord[i];
1128+
Vars->Mon2D_Buffer[i + ParticleCount*(Vars->Coord_Number+1)] = Coord[i];
11111129
}
11121130
#pragma acc atomic
11131131
Vars->Buffer_Counter = Vars->Buffer_Counter + 1;
11141132
if (Vars->Flag_Verbose && (Vars->Buffer_Counter >= Vars->Buffer_Block) && (Vars->Flag_List == 1))
11151133
printf("Monitor_nD: %s %li photons stored in List.\n", Vars->compcurname, Vars->Buffer_Counter);
11161134
}
1117-
#pragma acc atomic
1118-
Vars->Photon_Counter = Vars->Photon_Counter + 1;
11191135
} /* end (Vars->Flag_Auto_Limits != 2) */
11201136

11211137
} /* end while */

0 commit comments

Comments
 (0)