diff --git a/sixtracklib/block.c b/sixtracklib/block.c index 5dc632b7..512d184e 100644 --- a/sixtracklib/block.c +++ b/sixtracklib/block.c @@ -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, @@ -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, @@ -96,24 +143,23 @@ 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; @@ -121,15 +167,18 @@ CLKERNEL void Block_track(CLGLOBAL value_t *elems, 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 } diff --git a/sixtracklib/cblock.py b/sixtracklib/cblock.py index d9e90eba..6bdc7dca 100644 --- a/sixtracklib/cblock.py +++ b/sixtracklib/cblock.py @@ -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 @@ -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) @@ -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 @@ -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, diff --git a/sixtracklib/cobjects.py b/sixtracklib/cobjects.py index 2457f5ef..1c9dd544 100644 --- a/sixtracklib/cobjects.py +++ b/sixtracklib/cobjects.py @@ -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() diff --git a/sixtracklib/cparticles.py b/sixtracklib/cparticles.py index 14f0a9f0..5a2fbd04 100644 --- a/sixtracklib/cparticles.py +++ b/sixtracklib/cparticles.py @@ -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') diff --git a/sixtracklib/particle.h b/sixtracklib/particle.h index 7f52a1b8..e9dce0a5 100644 --- a/sixtracklib/particle.h +++ b/sixtracklib/particle.h @@ -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) ); diff --git a/sixtracklib/track.c b/sixtracklib/track.c index b0d47779..9776cdbd 100644 --- a/sixtracklib/track.c +++ b/sixtracklib/track.c @@ -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]; diff --git a/tests/test_track.py b/tests/test_track.py index 7f03e2a1..b9d31fa5 100644 --- a/tests/test_track.py +++ b/tests/test_track.py @@ -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