Skip to content

Commit 29f2028

Browse files
bernhardmgruberbernhardmgruber
committed
Refactoring
1 parent 8192121 commit 29f2028

File tree

6 files changed

+81
-75
lines changed

6 files changed

+81
-75
lines changed

‎examples/TestEm3/electrons.cu‎

Lines changed: 14 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -40,11 +40,14 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
4040

4141
int activeSize = active->size();
4242
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < activeSize; i += blockDim.x * gridDim.x) {
43-
const int slot = (*active)[i];
44-
Track &currentTrack = electrons[slot];
45-
auto volume = currentTrack.navState.Top();
46-
int volumeID = volume->id();
47-
int theMCIndex = MCIndex[volumeID];
43+
const int slot = (*active)[i];
44+
Track &currentTrack = electrons[slot];
45+
const int volumeID = currentTrack.navState.Top()->id();
46+
const int theMCIndex = MCIndex[volumeID];
47+
48+
auto survive = [&] {
49+
activeQueue->push_back(slot);
50+
};
4851

4952
// Init a track with the needed data to call into G4HepEm.
5053
G4HepEmElectronTrack elTrack;
@@ -218,21 +221,21 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
218221

219222
// Kill the particle if it left the world.
220223
if (nextState.Top() != nullptr) {
221-
activeQueue->push_back(slot);
222224
BVHNavigator::RelocateToNextVolume(currentTrack.pos, currentTrack.dir, nextState);
223225

224226
// Move to the next boundary.
225227
currentTrack.navState = nextState;
228+
survive();
226229
}
227230
continue;
228231
} else if (!propagated) {
229232
// Did not yet reach the interaction point due to error in the magnetic
230233
// field propagation. Try again next time.
231-
activeQueue->push_back(slot);
234+
survive();
232235
continue;
233236
} else if (winnerProcessIndex < 0) {
234237
// No discrete process, move on.
235-
activeQueue->push_back(slot);
238+
survive();
236239
continue;
237240
}
238241

@@ -243,7 +246,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
243246
// Check if a delta interaction happens instead of the real discrete process.
244247
if (G4HepEmElectronManager::CheckDelta(&g4HepEmData, theTrack, currentTrack.Uniform())) {
245248
// A delta interaction happened, move on.
246-
activeQueue->push_back(slot);
249+
survive();
247250
continue;
248251
}
249252

@@ -277,8 +280,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
277280

278281
currentTrack.energy = energy - deltaEkin;
279282
currentTrack.dir.Set(dirPrimary[0], dirPrimary[1], dirPrimary[2]);
280-
// The current track continues to live.
281-
activeQueue->push_back(slot);
283+
survive();
282284
break;
283285
}
284286
case 1: {
@@ -304,8 +306,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
304306

305307
currentTrack.energy = energy - deltaEkin;
306308
currentTrack.dir.Set(dirPrimary[0], dirPrimary[1], dirPrimary[2]);
307-
// The current track continues to live.
308-
activeQueue->push_back(slot);
309+
survive();
309310
break;
310311
}
311312
case 2: {

‎examples/TestEm3/gammas.cu‎

Lines changed: 13 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -25,11 +25,14 @@ __global__ void TransportGammas(Track *gammas, const adept::MParray *active, Sec
2525
{
2626
int activeSize = active->size();
2727
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < activeSize; i += blockDim.x * gridDim.x) {
28-
const int slot = (*active)[i];
29-
Track &currentTrack = gammas[slot];
30-
auto volume = currentTrack.navState.Top();
31-
int volumeID = volume->id();
32-
int theMCIndex = MCIndex[volumeID];
28+
const int slot = (*active)[i];
29+
Track &currentTrack = gammas[slot];
30+
const int volumeID = currentTrack.navState.Top()->id();
31+
const int theMCIndex = MCIndex[volumeID];
32+
33+
auto survive = [&] {
34+
activeQueue->push_back(slot);
35+
};
3336

3437
// Init a track with the needed data to call into G4HepEm.
3538
G4HepEmGammaTrack gammaTrack;
@@ -86,16 +89,16 @@ __global__ void TransportGammas(Track *gammas, const adept::MParray *active, Sec
8689

8790
// Kill the particle if it left the world.
8891
if (nextState.Top() != nullptr) {
89-
activeQueue->push_back(slot);
9092
BVHNavigator::RelocateToNextVolume(currentTrack.pos, currentTrack.dir, nextState);
9193

9294
// Move to the next boundary.
9395
currentTrack.navState = nextState;
96+
survive();
9497
}
9598
continue;
9699
} else if (winnerProcessIndex < 0) {
97100
// No discrete process, move on.
98-
activeQueue->push_back(slot);
101+
survive();
99102
continue;
100103
}
101104

@@ -114,7 +117,7 @@ __global__ void TransportGammas(Track *gammas, const adept::MParray *active, Sec
114117
case 0: {
115118
// Invoke gamma conversion to e-/e+ pairs, if the energy is above the threshold.
116119
if (energy < 2 * copcore::units::kElectronMassC2) {
117-
activeQueue->push_back(slot);
120+
survive();
118121
continue;
119122
}
120123

@@ -151,7 +154,7 @@ __global__ void TransportGammas(Track *gammas, const adept::MParray *active, Sec
151154
// Invoke Compton scattering of gamma.
152155
constexpr double LowEnergyThreshold = 100 * copcore::units::eV;
153156
if (energy < LowEnergyThreshold) {
154-
activeQueue->push_back(slot);
157+
survive();
155158
continue;
156159
}
157160
const double origDirPrimary[] = {currentTrack.dir.x(), currentTrack.dir.y(), currentTrack.dir.z()};
@@ -180,9 +183,7 @@ __global__ void TransportGammas(Track *gammas, const adept::MParray *active, Sec
180183
if (newEnergyGamma > LowEnergyThreshold) {
181184
currentTrack.energy = newEnergyGamma;
182185
currentTrack.dir = newDirGamma;
183-
184-
// The current track continues to live.
185-
activeQueue->push_back(slot);
186+
survive();
186187
} else {
187188
atomicAdd(&globalScoring->energyDeposit, newEnergyGamma);
188189
atomicAdd(&scoringPerVolume->energyDeposit[volumeID], newEnergyGamma);

‎examples/TestEm3Compact/electrons.cu‎

Lines changed: 14 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -39,11 +39,14 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
3939

4040
int activeSize = electrons->fActiveTracks->size();
4141
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < activeSize; i += blockDim.x * gridDim.x) {
42-
const int slot = (*electrons->fActiveTracks)[i];
43-
Track &currentTrack = (*electrons)[slot];
44-
auto volume = currentTrack.navState.Top();
45-
int volumeID = volume->id();
46-
int theMCIndex = MCIndex[volumeID];
42+
const int slot = (*electrons->fActiveTracks)[i];
43+
Track &currentTrack = (*electrons)[slot];
44+
const int volumeID = currentTrack.navState.Top()->id();
45+
const int theMCIndex = MCIndex[volumeID];
46+
47+
auto survive = [&] {
48+
electrons->fNextTracks->push_back(slot);
49+
};
4750

4851
// Init a track with the needed data to call into G4HepEm.
4952
G4HepEmElectronTrack elTrack;
@@ -217,21 +220,21 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
217220

218221
// Kill the particle if it left the world.
219222
if (nextState.Top() != nullptr) {
220-
electrons->fNextTracks->push_back(slot);
221223
BVHNavigator::RelocateToNextVolume(currentTrack.pos, currentTrack.dir, nextState);
222224

223225
// Move to the next boundary.
224226
currentTrack.navState = nextState;
227+
survive();
225228
}
226229
continue;
227230
} else if (!propagated) {
228231
// Did not yet reach the interaction point due to error in the magnetic
229232
// field propagation. Try again next time.
230-
electrons->fNextTracks->push_back(slot);
233+
survive();
231234
continue;
232235
} else if (winnerProcessIndex < 0) {
233236
// No discrete process, move on.
234-
electrons->fNextTracks->push_back(slot);
237+
survive();
235238
continue;
236239
}
237240

@@ -242,7 +245,7 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
242245
// Check if a delta interaction happens instead of the real discrete process.
243246
if (G4HepEmElectronManager::CheckDelta(&g4HepEmData, theTrack, currentTrack.Uniform())) {
244247
// A delta interaction happened, move on.
245-
electrons->fNextTracks->push_back(slot);
248+
survive();
246249
continue;
247250
}
248251

@@ -276,8 +279,7 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
276279

277280
currentTrack.energy = energy - deltaEkin;
278281
currentTrack.dir.Set(dirPrimary[0], dirPrimary[1], dirPrimary[2]);
279-
// The current track continues to live.
280-
electrons->fNextTracks->push_back(slot);
282+
survive();
281283
break;
282284
}
283285
case 1: {
@@ -303,8 +305,7 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
303305

304306
currentTrack.energy = energy - deltaEkin;
305307
currentTrack.dir.Set(dirPrimary[0], dirPrimary[1], dirPrimary[2]);
306-
// The current track continues to live.
307-
electrons->fNextTracks->push_back(slot);
308+
survive();
308309
break;
309310
}
310311
case 2: {

‎examples/TestEm3Compact/gammas.cu‎

Lines changed: 13 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -24,11 +24,14 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
2424
{
2525
int activeSize = gammas->fActiveTracks->size();
2626
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < activeSize; i += blockDim.x * gridDim.x) {
27-
const int slot = (*gammas->fActiveTracks)[i];
28-
Track &currentTrack = (*gammas)[slot];
29-
auto volume = currentTrack.navState.Top();
30-
int volumeID = volume->id();
31-
int theMCIndex = MCIndex[volumeID];
27+
const int slot = (*gammas->fActiveTracks)[i];
28+
Track &currentTrack = (*gammas)[slot];
29+
const int volumeID = currentTrack.navState.Top()->id();
30+
const int theMCIndex = MCIndex[volumeID];
31+
32+
auto survive = [&] {
33+
gammas->fNextTracks->push_back(slot);
34+
};
3235

3336
// Init a track with the needed data to call into G4HepEm.
3437
G4HepEmGammaTrack gammaTrack;
@@ -85,16 +88,16 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
8588

8689
// Kill the particle if it left the world.
8790
if (nextState.Top() != nullptr) {
88-
gammas->fNextTracks->push_back(slot);
8991
BVHNavigator::RelocateToNextVolume(currentTrack.pos, currentTrack.dir, nextState);
9092

9193
// Move to the next boundary.
9294
currentTrack.navState = nextState;
95+
survive();
9396
}
9497
continue;
9598
} else if (winnerProcessIndex < 0) {
9699
// No discrete process, move on.
97-
gammas->fNextTracks->push_back(slot);
100+
survive();
98101
continue;
99102
}
100103

@@ -113,7 +116,7 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
113116
case 0: {
114117
// Invoke gamma conversion to e-/e+ pairs, if the energy is above the threshold.
115118
if (energy < 2 * copcore::units::kElectronMassC2) {
116-
gammas->fNextTracks->push_back(slot);
119+
survive();
117120
continue;
118121
}
119122

@@ -150,7 +153,7 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
150153
// Invoke Compton scattering of gamma.
151154
constexpr double LowEnergyThreshold = 100 * copcore::units::eV;
152155
if (energy < LowEnergyThreshold) {
153-
gammas->fNextTracks->push_back(slot);
156+
survive();
154157
continue;
155158
}
156159
const double origDirPrimary[] = {currentTrack.dir.x(), currentTrack.dir.y(), currentTrack.dir.z()};
@@ -179,9 +182,7 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
179182
if (newEnergyGamma > LowEnergyThreshold) {
180183
currentTrack.energy = newEnergyGamma;
181184
currentTrack.dir = newDirGamma;
182-
183-
// The current track continues to live.
184-
gammas->fNextTracks->push_back(slot);
185+
survive();
185186
} else {
186187
atomicAdd(&globalScoring->energyDeposit, newEnergyGamma);
187188
atomicAdd(&scoringPerVolume->energyDeposit[volumeID], newEnergyGamma);

‎examples/TestEm3MT/electrons.cu‎

Lines changed: 14 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -40,11 +40,14 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
4040

4141
int activeSize = active->size();
4242
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < activeSize; i += blockDim.x * gridDim.x) {
43-
const int slot = (*active)[i];
44-
Track &currentTrack = electrons[slot];
45-
auto volume = currentTrack.navState.Top();
46-
int volumeID = volume->id();
47-
int theMCIndex = MCIndex[volumeID];
43+
const int slot = (*active)[i];
44+
Track &currentTrack = electrons[slot];
45+
const int volumeID = currentTrack.navState.Top()->id();
46+
const int theMCIndex = MCIndex[volumeID];
47+
48+
auto survive = [&] {
49+
activeQueue->push_back(slot);
50+
};
4851

4952
// Init a track with the needed data to call into G4HepEm.
5053
G4HepEmElectronTrack elTrack;
@@ -219,21 +222,21 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
219222

220223
// Kill the particle if it left the world.
221224
if (nextState.Top() != nullptr) {
222-
activeQueue->push_back(slot);
223225
BVHNavigator::RelocateToNextVolume(currentTrack.pos, currentTrack.dir, nextState);
224226

225227
// Move to the next boundary.
226228
currentTrack.navState = nextState;
229+
survive();
227230
}
228231
continue;
229232
} else if (!propagated) {
230233
// Did not yet reach the interaction point due to error in the magnetic
231234
// field propagation. Try again next time.
232-
activeQueue->push_back(slot);
235+
survive();
233236
continue;
234237
} else if (winnerProcessIndex < 0) {
235238
// No discrete process, move on.
236-
activeQueue->push_back(slot);
239+
survive();
237240
continue;
238241
}
239242

@@ -244,7 +247,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
244247
// Check if a delta interaction happens instead of the real discrete process.
245248
if (G4HepEmElectronManager::CheckDelta(&g4HepEmData, theTrack, currentTrack.Uniform())) {
246249
// A delta interaction happened, move on.
247-
activeQueue->push_back(slot);
250+
survive();
248251
continue;
249252
}
250253

@@ -278,8 +281,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
278281

279282
currentTrack.energy = energy - deltaEkin;
280283
currentTrack.dir.Set(dirPrimary[0], dirPrimary[1], dirPrimary[2]);
281-
// The current track continues to live.
282-
activeQueue->push_back(slot);
284+
survive();
283285
break;
284286
}
285287
case 1: {
@@ -305,8 +307,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
305307

306308
currentTrack.energy = energy - deltaEkin;
307309
currentTrack.dir.Set(dirPrimary[0], dirPrimary[1], dirPrimary[2]);
308-
// The current track continues to live.
309-
activeQueue->push_back(slot);
310+
survive();
310311
break;
311312
}
312313
case 2: {

0 commit comments

Comments
 (0)