Skip to content

Commit

Permalink
Merge pull request #5 from rdemaria/2018_reload
Browse files Browse the repository at this point in the history 
2018 reload
  • Loading branch information
@giadarol
giadarol authored Feb 20, 2018
2 parents 9ba4720 + 475b2d0 commit e7f7b55
Show file tree
Hide file tree
Showing 7 changed files with 94 additions and 25 deletions.
73 changes: 61 additions & 12 deletions sixtracklib/block.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -44,10 +44,13 @@ double Align_get_dy(CLGLOBAL value_t *elem){return elem[4].f64;}
#ifdef _GPUCODE


void track_single(Particles *particles, uint64_t partid, CLGLOBAL value_t * elem){
void track_single(Particles *particles, uint64_t partid,
CLGLOBAL value_t * elem){
//printf("single: partid=%u, typeid=%zu\n",partid, elem[0].u64);
if (particles->state[partid] >= 0 ) {
//printf("single: typeid=%zu\n",elem[0].u64);
enum type_t typeid = get_type(elem);
// _DP("Block_track: elemid=%zu typedid=%u\n",elemid,typeid);
//printf("single: typedid=%u\n",typeid);
switch (typeid) {
case DriftID:
Drift_track(particles, partid,
Expand Down Expand Up @@ -82,6 +85,50 @@ void track_single(Particles *particles, uint64_t partid, CLGLOBAL value_t * elem
}//end if state
}


CLKERNEL void Block_unpack(
CLGLOBAL value_t *particles_p, //Particles
CLGLOBAL value_t *elembyelem_p, //ElembyElem
CLGLOBAL value_t *turnbyturn_p) //TurnbyTurn
{
uint64_t partid = get_global_id(0);

bool elembyelem_flag = (elembyelem_p[0].i64 != 0);
bool turnbyturn_flag = (turnbyturn_p[0].i64 != 0);

Particles* particles = (Particles*) particles_p;
Particles* elembyelem = (Particles*) elembyelem_p;
Particles* turnbyturn = (Particles*) turnbyturn_p;

Particles_unpack( particles );

//printf( "p->beta0[0] %g\n", particles->beta0[0]);
//printf( "p->beta0[3] %g\n", particles->beta0[3]);
//printf( "elembyelem[0] %i \n", elembyelem_p[0].u64);
//printf( "turnbyturn[0] %i \n", turnbyturn_p[0].u64);
//printf( "elembyelem[1] %i \n", elembyelem_p[1].u64);
//printf( "turnbyturn[1] %i \n", turnbyturn_p[1].u64);
//printf( "elembyelem[2] %i \n", elembyelem_p[2].u64);
//printf( "turnbyturn[2] %i \n", turnbyturn_p[2].u64);


if (elembyelem_flag) Particles_unpack( elembyelem );
if (turnbyturn_flag) Particles_unpack( turnbyturn );

//printf( "elembyelem[0] %i \n", elembyelem_p[0].u64);
//printf( "turnbyturn[0] %i \n", turnbyturn_p[0].u64);
//printf( "elembyelem[1] %i \n", elembyelem_p[1].u64);
//printf( "turnbyturn[1] %i \n", turnbyturn_p[1].u64);
//printf( "elembyelem[2] %i \n", elembyelem_p[2].u64);
//printf( "turnbyturn[2] %i \n", turnbyturn_p[2].u64);

//printf( "ele->beta0[0] %g\n", elembyelem->beta0[0]);
//printf( "ele->beta0[3] %g\n", elembyelem->beta0[3]);
};




CLKERNEL void Block_track(CLGLOBAL value_t *elems,
CLGLOBAL uint64_t *elemids,
uint64_t nelems,
Expand All @@ -96,40 +143,42 @@ CLKERNEL void Block_track(CLGLOBAL value_t *elems,

Particles* particles = (Particles*) particles_p;

printf( "%s test\n","test" );
//printf( "beta0[%d] %g\n",partid, particles->beta0[partid]);

bool elembyelem_flag = (elembyelem_p[0].i64 >= 0);
bool turnbyturn_flag = (turnbyturn_p[0].i64 >= 0);
bool elembyelem_flag = (elembyelem_p[0].i64 != 0);
bool turnbyturn_flag = (turnbyturn_p[0].i64 != 0);

Particles* elembyelem = (Particles*) elembyelem_p;
Particles* turnbyturn = (Particles*) turnbyturn_p;

Particles_unpack( particles);
if (elembyelem_flag) {
if (partid==0) Particles_unpack( elembyelem );
Particles_copy(particles, elembyelem, partid, partid);
};
if (turnbyturn_flag) {
if (partid==0) Particles_unpack( turnbyturn );
Particles_copy(particles, turnbyturn, partid, partid);
};

//printf( "tbt->beta0[%d] %g\n",partid, turnbyturn->beta0[partid]);

uint64_t nparts=particles->npart;
uint64_t tbt=nparts;
uint64_t ebe=nparts;

for (int jj = 0; jj < nturns; jj++) {
for (int ii = 0; ii < nelems; ii++) {
elemid = elemids[ii];
//printf("elemid %u\n",elemid);
elem = elems+elemid;
track_single(particles,partid,elem);
if (elembyelem_flag)
ebe+=nparts;
if (elembyelem_flag){
Particles_copy(particles, elembyelem, partid, ebe+partid);
ebe+=nparts;
}
} //end elem loop
if (turnbyturn_flag)
tbt+=nparts;
if (turnbyturn_flag){
Particles_copy(particles, turnbyturn, partid, tbt+partid);
tbt+=nparts;
}
} //end turn loop
}

Expand Down
14 changes: 8 additions & 6 deletions sixtracklib/cblock.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
os.environ['PYOPENCL_COMPILER_OUTPUT']='1'
srcpath='-I%s'%modulepath
src=open(os.path.join(modulepath,'block.c')).read()
ctx = cl.create_some_context(interactive=False)
ctx = cl.create_some_context(interactive=False,answers=[1])
prg=cl.Program(ctx,src).build(options=[srcpath])
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags
Expand All @@ -36,7 +36,7 @@ class DriftExact(CObject):

class Multipole(CObject):
objid = CProp('u64',0,default=4)
order = CProp('f64',1,default=0,const=True)
order = CProp('u64',1,default=0,const=True)
length = CProp('f64',2,default=0)
hxl = CProp('f64',3,default=0)
hyl = CProp('f64',4,default=0)
Expand Down Expand Up @@ -96,16 +96,16 @@ def track_cl(self,particles,nturns=1,elembyelem=None,turnbyturn=None):
particles_g=cl.Buffer(ctx, rw, hostbuf=particles._data)
#ElemByElem data
if elembyelem is True:
elembyelem=CParticles(npart=npart*self.nelems*nturns+1)
elembyelem=CParticles(npart=npart*(self.nelems*nturns+1))
if elembyelem is None:
elembyelem_g=cl.Buffer(ctx, rw, hostbuf=np.array([-1]))
elembyelem_g=cl.Buffer(ctx, rw, hostbuf=np.array([0]))
else:
elembyelem_g=cl.Buffer(ctx, rw, hostbuf=elembyelem._data)
#TurnByTurn data
if turnbyturn is True:
turnbyturn=CParticles(npart=npart*nturns+1)
turnbyturn=CParticles(npart=npart*(nturns+1))
if turnbyturn is None:
turnbyturn_g=cl.Buffer(ctx, rw, hostbuf=np.array([-1]))
turnbyturn_g=cl.Buffer(ctx, rw, hostbuf=np.array([0]))
else:
turnbyturn_g=cl.Buffer(ctx, rw, hostbuf=turnbyturn._data)
#Tracking data
Expand All @@ -114,6 +114,8 @@ def track_cl(self,particles,nturns=1,elembyelem=None,turnbyturn=None):
elemids_g=cl.Buffer(ctx, rw, hostbuf=elemids)
nelems=np.int64(self.nelems)
nturns=np.int64(nturns)
prg.Block_unpack(queue,[1],None,
particles_g, elembyelem_g, turnbyturn_g)
prg.Block_track(queue,[npart],None,
elems_g, elemids_g, nelems,
nturns,
Expand Down
11 changes: 10 additions & 1 deletion sixtracklib/cobjects.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -134,4 +134,13 @@ def _fill_args(self,nvargs):
nvargs.get(name,prop.default)
elif isinstance(prop.valuetype,CObject):
raise NotImplemented('Nested object not implemented')

def pretty_print(self):
props=self._get_props()
out=['%s('%(self.__class__.__name__)]
for offset,name,prop in props:
val='%s'%getattr(self,name)
out.append(' %-8s = %s,'%(name,val[:70]))
out.append(')')
return '\n'.join(out)
def __repr__(self):
return self.pretty_print()
8 changes: 4 additions & 4 deletions sixtracklib/cparticles.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,11 +17,11 @@ class CParticles(CObject):
kboltz=8.6173303e-5 #ev K^-1 #1.38064852e-23 # JK^-1

npart =CProp('u64', 0, 0, const=True)
q0 =CProp('f64', 1, 1, length='npart' )
q0 =CProp('f64', 1, 1, length='npart')
mass0 =CProp('f64', 2, pmass, length='npart')
p0c =CProp('f64', 5, 450e6, length='npart')
beta0 =CProp('f64', 3, 450e6/np.sqrt(450e6**2+pmass**2), length='npart')
gamma0=CProp('f64', 4, np.sqrt(450e6**2+pmass**2)/pmass, length='npart')
p0c =CProp('f64', 5, 450e9, length='npart')
beta0 =CProp('f64', 3, 450e9/np.sqrt(450e9**2+pmass**2), length='npart')
gamma0=CProp('f64', 4, np.sqrt(450e9**2+pmass**2)/pmass, length='npart')
partid=CProp('u64', 6, 0, length='npart')
elemid=CProp('u64', 7, 0, length='npart')
turn =CProp('u64', 8, 0, length='npart')
Expand Down
5 changes: 5 additions & 0 deletions sixtracklib/particle.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,11 @@ typedef CLGLOBAL struct Particles {
} Particles;

Particles* Particles_unpack(Particles* p) {
p->q0 = ( (CLGLOBAL double *) p + ((uint64_t) p->q0 ) );
p->mass0 = ( (CLGLOBAL double *) p + ((uint64_t) p->mass0 ) );
p->beta0 = ( (CLGLOBAL double *) p + ((uint64_t) p->beta0 ) );
p->gamma0 = ( (CLGLOBAL double *) p + ((uint64_t) p->gamma0) );
p->p0c = ( (CLGLOBAL double *) p + ((uint64_t) p->p0c ) );
p->partid = ( (CLGLOBAL int64_t *) p + ((uint64_t) p->partid) );
p->elemid = ( (CLGLOBAL int64_t *) p + ((uint64_t) p->elemid) );
p->turn = ( (CLGLOBAL int64_t *) p + ((uint64_t) p->turn) );
Expand Down
5 changes: 4 additions & 1 deletion sixtracklib/track.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -57,10 +57,13 @@ inline int DriftExact_track(Particles* p, uint64_t ip,


inline int Multipole_track(Particles* p, uint64_t ip,
long int order, double l,
uint64_t order, double l,
double hxl, double hyl, CLGLOBAL double* bal){
double x,y,chi,dpx,dpy,zre,zim,b1l,a1l,hxx,hyy;

//printf("multipole: partid:%d, order:%d, bal[2]:%g \n",
// ip,order,bal[2]);

dpx = bal[order*2];
dpy = bal[order*2+1];
x = p->x[ip]; y = p->y[ip]; chi = p->chi[ip];
Expand Down
3 changes: 2 additions & 1 deletion tests/test_track.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,8 @@ def test_track():
bunch.y[2]=0.2
bunch.sigma[3]=0.1

particles,ebe,tbt=fodo.track_cl(bunch,1,True,True)
particles,ebe,tbt=fodo.track_cl(bunch,nturns=1,
elembyelem=True,turnbyturn=True)
return particles,ebe,tbt


Expand Down

0 comments on commit e7f7b55

Please sign in to comment.