Skip to content

Commit a814ad5

Browse files
psychocoderHPCpsychocoderHPC
committed
[examples] transform kernel functions to functors
- use `PMACC_TYPEKERNEL` macro to start a kernel - transform global kernel function to a functor
1 parent d0aa883 commit a814ad5

File tree

3 files changed

+159
-153
lines changed

3 files changed

+159
-153
lines changed

‎examples/SingleParticleCurrent/include/particles/ParticlesInitOneParticle.hpp

Lines changed: 52 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -34,67 +34,69 @@
3434
namespace picongpu
3535
{
3636

37-
template< class ParBox>
38-
__global__ void kernelAddOneParticle(ParBox pb,
39-
DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell)
37+
struct kernelAddOneParticle
4038
{
41-
typedef typename ParBox::FramePtr FramePtr;
42-
43-
FramePtr frame;
44-
45-
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
39+
template< class ParBox >
40+
DINLINE void operator()(ParBox pb, DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell) const
41+
{
42+
typedef typename ParBox::FramePtr FramePtr;
4643

47-
float_X parWeighting = particles::TYPICAL_NUM_PARTICLES_PER_MACROPARTICLE;
44+
FramePtr frame;
4845

49-
frame = pb.getEmptyFrame();
50-
pb.setAsLastFrame(frame, superCell);
46+
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
5147

48+
float_X parWeighting = particles::TYPICAL_NUM_PARTICLES_PER_MACROPARTICLE;
5249

50+
frame = pb.getEmptyFrame();
51+
pb.setAsLastFrame(frame, superCell);
5352

5453

55-
// many particle loop:
56-
for (unsigned i = 0; i < 1; ++i)
57-
{
5854

59-
PMACC_AUTO(par, frame[i]);
6055

61-
/** we now initialize all attributes of the new particle to their default values
62-
* some attributes, such as the position, localCellIdx, weighting or the
63-
* multiMask (\see AttrToIgnore) of the particle will be set individually
64-
* in the following lines since they are already known at this point.
65-
*/
56+
// many particle loop:
57+
for (unsigned i = 0; i < 1; ++i)
6658
{
67-
typedef typename ParBox::FrameType FrameType;
68-
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
69-
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
70-
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
71-
72-
algorithms::forEach::ForEach<ParticleCleanedAttrList,
73-
SetAttributeToDefault<bmpl::_1> > setToDefault;
74-
setToDefault(forward(par));
59+
60+
PMACC_AUTO(par, frame[i]);
61+
62+
/** we now initialize all attributes of the new particle to their default values
63+
* some attributes, such as the position, localCellIdx, weighting or the
64+
* multiMask (\see AttrToIgnore) of the particle will be set individually
65+
* in the following lines since they are already known at this point.
66+
*/
67+
{
68+
typedef typename ParBox::FrameType FrameType;
69+
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
70+
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
71+
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
72+
73+
algorithms::forEach::ForEach<ParticleCleanedAttrList,
74+
SetAttributeToDefault<bmpl::_1> > setToDefault;
75+
setToDefault(forward(par));
76+
}
77+
floatD_X pos = float3_X(LOCAL_POS_X, LOCAL_POS_Y, LOCAL_POS_Z).shrink<simDim>();
78+
79+
const float_X GAMMA0 = (float_X) (1.0 / sqrt(1.0 - (BETA0_X * BETA0_X + BETA0_Y * BETA0_Y + BETA0_Z * BETA0_Z)));
80+
float3_X mom = float3_X(
81+
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
82+
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
83+
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
84+
);
85+
86+
par[position_] = pos;
87+
par[momentum_] = mom;
88+
par[multiMask_] = 1;
89+
par[localCellIdx_] = linearIdx;
90+
par[weighting_] = parWeighting;
91+
92+
#if(ENABLE_RADIATION == 1)
93+
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
94+
par[radiationFlag_] = true;
95+
#endif
96+
#endif
7597
}
76-
floatD_X pos = float3_X(LOCAL_POS_X, LOCAL_POS_Y, LOCAL_POS_Z).shrink<simDim>();
77-
78-
const float_X GAMMA0 = (float_X) (1.0 / sqrt(1.0 - (BETA0_X * BETA0_X + BETA0_Y * BETA0_Y + BETA0_Z * BETA0_Z)));
79-
float3_X mom = float3_X(
80-
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
81-
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
82-
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
83-
);
84-
85-
par[position_] = pos;
86-
par[momentum_] = mom;
87-
par[multiMask_] = 1;
88-
par[localCellIdx_] = linearIdx;
89-
par[weighting_] = parWeighting;
90-
91-
#if(ENABLE_RADIATION == 1)
92-
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
93-
par[radiationFlag_] = true;
94-
#endif
95-
#endif
9698
}
97-
}
99+
};
98100

99101
template<class ParticlesClass>
100102
class ParticlesInitOneParticle
@@ -128,7 +130,7 @@ class ParticlesInitOneParticle
128130
std::cout << "localSuperCell: " << localSuperCell << std::endl;
129131
std::cout << "cellInSuperCell: " << cellInSuperCell << std::endl;
130132

131-
__cudaKernel(kernelAddOneParticle)
133+
PMACC_TYPEKERNEL(kernelAddOneParticle)
132134
(1, 1)
133135
(parClass.getDeviceParticlesBox(),
134136
localSuperCell, cellInSuperCell);

‎examples/SingleParticleRadiationWithLaser/include/particles/ParticlesInitOneParticle.hpp

Lines changed: 54 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -37,69 +37,71 @@
3737
namespace picongpu
3838
{
3939

40-
template< class ParBox>
41-
__global__ void kernelAddOneParticle(ParBox pb,
42-
DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell)
40+
struct kernelAddOneParticle
4341
{
44-
typedef typename ParBox::FramePtr FramePtr;
45-
46-
FramePtr frame;
42+
template< class ParBox>
43+
DINLINE void operator()(ParBox pb, DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell) const
44+
{
45+
typedef typename ParBox::FramePtr FramePtr;
4746

48-
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
47+
FramePtr frame;
4948

50-
float_X parWeighting = particles::TYPICAL_NUM_PARTICLES_PER_MACROPARTICLE;
49+
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
5150

52-
frame = pb.getEmptyFrame();
53-
pb.setAsLastFrame(frame, superCell);
51+
float_X parWeighting = particles::TYPICAL_NUM_PARTICLES_PER_MACROPARTICLE;
5452

53+
frame = pb.getEmptyFrame();
54+
pb.setAsLastFrame(frame, superCell);
5555

5656

5757

58-
// many particle loop:
59-
for (unsigned i = 0; i < 1; ++i)
60-
{
61-
PMACC_AUTO(par, frame[i]);
6258

63-
/** we now initialize all attributes of the new particle to their default values
64-
* some attributes, such as the position, localCellIdx, weighting or the
65-
* multiMask (\see AttrToIgnore) of the particle will be set individually
66-
* in the following lines since they are already known at this point.
67-
*/
59+
// many particle loop:
60+
for (unsigned i = 0; i < 1; ++i)
6861
{
69-
typedef typename ParBox::FrameType FrameType;
70-
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
71-
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
72-
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
73-
74-
algorithms::forEach::ForEach<ParticleCleanedAttrList,
75-
SetAttributeToDefault<bmpl::_1> > setToDefault;
76-
setToDefault(forward(par));
62+
PMACC_AUTO(par, frame[i]);
63+
64+
/** we now initialize all attributes of the new particle to their default values
65+
* some attributes, such as the position, localCellIdx, weighting or the
66+
* multiMask (\see AttrToIgnore) of the particle will be set individually
67+
* in the following lines since they are already known at this point.
68+
*/
69+
{
70+
typedef typename ParBox::FrameType FrameType;
71+
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
72+
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
73+
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
74+
75+
algorithms::forEach::ForEach<ParticleCleanedAttrList,
76+
SetAttributeToDefault<bmpl::_1> > setToDefault;
77+
setToDefault(forward(par));
78+
}
79+
80+
floatD_X pos(floatD_X::create(0.5));
81+
82+
const float_X GAMMA0_X = 1.0f / sqrtf(1.0f - float_X(BETA0_X * BETA0_X));
83+
const float_X GAMMA0_Y = 1.0f / sqrtf(1.0f - float_X(BETA0_Y * BETA0_Y));
84+
const float_X GAMMA0_Z = 1.0f / sqrtf(1.0f - float_X(BETA0_Z * BETA0_Z));
85+
float3_X mom = float3_X(
86+
GAMMA0_X * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
87+
GAMMA0_Y * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
88+
GAMMA0_Z * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
89+
);
90+
91+
par[position_] = pos;
92+
par[momentum_] = mom;
93+
par[multiMask_] = 1;
94+
par[localCellIdx_] = linearIdx;
95+
par[weighting_] = parWeighting;
96+
97+
#if(ENABLE_RADIATION == 1)
98+
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
99+
par[radiationFlag_] = true;
100+
#endif
101+
#endif
77102
}
78-
79-
floatD_X pos(floatD_X::create(0.5));
80-
81-
const float_X GAMMA0_X = 1.0f / sqrtf(1.0f - float_X(BETA0_X * BETA0_X));
82-
const float_X GAMMA0_Y = 1.0f / sqrtf(1.0f - float_X(BETA0_Y * BETA0_Y));
83-
const float_X GAMMA0_Z = 1.0f / sqrtf(1.0f - float_X(BETA0_Z * BETA0_Z));
84-
float3_X mom = float3_X(
85-
GAMMA0_X * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
86-
GAMMA0_Y * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
87-
GAMMA0_Z * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
88-
);
89-
90-
par[position_] = pos;
91-
par[momentum_] = mom;
92-
par[multiMask_] = 1;
93-
par[localCellIdx_] = linearIdx;
94-
par[weighting_] = parWeighting;
95-
96-
#if(ENABLE_RADIATION == 1)
97-
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
98-
par[radiationFlag_] = true;
99-
#endif
100-
#endif
101103
}
102-
}
104+
};
103105

104106
template<class ParticlesClass>
105107
class ParticlesInitOneParticle
@@ -129,7 +131,7 @@ class ParticlesInitOneParticle
129131
localSuperCell = localSuperCell + cellDescription.getGuardingSuperCells();
130132

131133

132-
__cudaKernel(kernelAddOneParticle)
134+
PMACC_TYPEKERNEL(kernelAddOneParticle)
133135
(1, 1)
134136
(parClass.getDeviceParticlesBox(),
135137
localSuperCell, cellInSuperCell);

‎examples/SingleParticleTest/include/particles/ParticlesInitOneParticle.hpp

Lines changed: 53 additions & 51 deletions
Original file line numberDiff line numberDiff line change
@@ -34,68 +34,70 @@
3434
namespace picongpu
3535
{
3636

37-
template< class ParBox>
38-
__global__ void kernelAddOneParticle(ParBox pb,
39-
DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell)
37+
struct kernelAddOneParticle
4038
{
41-
typedef typename ParBox::FramePtr FramePtr;
42-
43-
FramePtr frame;
39+
template< class ParBox>
40+
DINLINE void operator()(ParBox pb, DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell) const
41+
{
42+
typedef typename ParBox::FramePtr FramePtr;
4443

45-
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
44+
FramePtr frame;
4645

47-
float_X parWeighting = particles::TYPICAL_NUM_PARTICLES_PER_MACROPARTICLE;
46+
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
4847

49-
frame = pb.getEmptyFrame();
50-
pb.setAsLastFrame(frame, superCell);
48+
float_X parWeighting = particles::TYPICAL_NUM_PARTICLES_PER_MACROPARTICLE;
5149

50+
frame = pb.getEmptyFrame();
51+
pb.setAsLastFrame(frame, superCell);
5252

5353

5454

55-
// many particle loop:
56-
for (unsigned i = 0; i < 1; ++i)
57-
{
58-
PMACC_AUTO(par, frame[i]);
5955

60-
/** we now initialize all attributes of the new particle to their default values
61-
* some attributes, such as the position, localCellIdx, weighting or the
62-
* multiMask (\see AttrToIgnore) of the particle will be set individually
63-
* in the following lines since they are already known at this point.
64-
*/
56+
// many particle loop:
57+
for (unsigned i = 0; i < 1; ++i)
6558
{
66-
typedef typename ParBox::FrameType FrameType;
67-
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
68-
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
69-
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
70-
71-
algorithms::forEach::ForEach<ParticleCleanedAttrList,
72-
SetAttributeToDefault<bmpl::_1> > setToDefault;
73-
setToDefault(forward(par));
59+
PMACC_AUTO(par, frame[i]);
60+
61+
/** we now initialize all attributes of the new particle to their default values
62+
* some attributes, such as the position, localCellIdx, weighting or the
63+
* multiMask (\see AttrToIgnore) of the particle will be set individually
64+
* in the following lines since they are already known at this point.
65+
*/
66+
{
67+
typedef typename ParBox::FrameType FrameType;
68+
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
69+
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
70+
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
71+
72+
algorithms::forEach::ForEach<ParticleCleanedAttrList,
73+
SetAttributeToDefault<bmpl::_1> > setToDefault;
74+
setToDefault(forward(par));
75+
}
76+
77+
floatD_X pos(floatD_X::create(0.5));
78+
79+
const float_X GAMMA0 = (float_X) (1.0 / sqrt(1.0 - (BETA0_X * BETA0_X + BETA0_Y * BETA0_Y + BETA0_Z * BETA0_Z)));
80+
float3_X mom = float3_X(
81+
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
82+
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
83+
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
84+
);
85+
86+
87+
par[position_] = pos;
88+
par[momentum_] = mom;
89+
par[multiMask_] = 1;
90+
par[localCellIdx_] = linearIdx;
91+
par[weighting_] = parWeighting;
92+
93+
#if(ENABLE_RADIATION == 1)
94+
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
95+
par[radiationFlag_] = true;
96+
#endif
97+
#endif
7498
}
75-
76-
floatD_X pos(floatD_X::create(0.5));
77-
78-
const float_X GAMMA0 = (float_X) (1.0 / sqrt(1.0 - (BETA0_X * BETA0_X + BETA0_Y * BETA0_Y + BETA0_Z * BETA0_Z)));
79-
float3_X mom = float3_X(
80-
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
81-
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
82-
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
83-
);
84-
85-
86-
par[position_] = pos;
87-
par[momentum_] = mom;
88-
par[multiMask_] = 1;
89-
par[localCellIdx_] = linearIdx;
90-
par[weighting_] = parWeighting;
91-
92-
#if(ENABLE_RADIATION == 1)
93-
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
94-
par[radiationFlag_] = true;
95-
#endif
96-
#endif
9799
}
98-
}
100+
};
99101

100102
template<class ParticlesClass>
101103
class ParticlesInitOneParticle
@@ -125,7 +127,7 @@ class ParticlesInitOneParticle
125127
localSuperCell = localSuperCell + cellDescription.getGuardingSuperCells();
126128

127129

128-
__cudaKernel(kernelAddOneParticle)
130+
PMACC_TYPEKERNEL(kernelAddOneParticle)
129131
(1, 1)
130132
(parClass.getDeviceParticlesBox(),
131133
localSuperCell, cellInSuperCell);

0 commit comments

Comments
 (0)