39 GPU_DATA_IS_MAPPED(h[0:p->
n_mrk])
40 GPU_PARALLEL_LOOP_ALL_LEVELS
41 for(i = 0; i < p->
n_mrk; i++) {
45 real k1[6], k2[6], k3[6], k4[6];
65 yprev[3] = p->
ppar[i];
67 yprev[5] = p->
zeta[i];
75 B_dB[4] = p->
B_phi[i];
90 step_gceom(k1, yprev, mass, charge, B_dB, E, aldforce);
96 for(
int j = 0; j < 6; j++) {
97 tempy[j] = yprev[j] + h[i]*k1[j]/2.0;
103 t0 + h[i]/2.0, Bdata);
107 t0 + h[i]/2.0, Edata, Bdata);
110 step_gceom(k2, tempy, mass, charge, B_dB, E, aldforce);
112 for(
int j = 0; j < 6; j++) {
113 tempy[j] = yprev[j] + h[i]*k2[j]/2.0;
119 t0 + h[i]/2.0, Bdata);
123 t0 + h[i]/2.0, Edata, Bdata);
126 step_gceom(k3, tempy, mass, charge, B_dB, E, aldforce);
128 for(
int j = 0; j < 6; j++) {
129 tempy[j] = yprev[j] + h[i]*k3[j];
139 t0 + h[i], Edata, Bdata);}
141 step_gceom(k4, tempy, mass, charge, B_dB, E, aldforce);
143 for(
int j = 0; j < 6; j++) {
145 + h[i]/6.0 * (k1[j] + 2*k2[j] + 2*k3[j] + k4[j]);
150 if(!errflag && y[0] <= 0) {
153 if(!errflag && y[4] < 0) {
191 p->
B_phi[i] = B_dB[4];
206 p->
theta[i] += atan2( (R0-axisrz[0]) * (p->
z[i]-axisrz[1])
207 - (z0-axisrz[1]) * (p->
r[i]-axisrz[0]),
208 (R0-axisrz[0]) * (p->
r[i]-axisrz[0])
209 + (z0-axisrz[1]) * (p->
z[i]-axisrz[1]) );
240 #pragma omp simd aligned(h : 64)
241 for(i = 0; i <
NSIMD; i++) {
245 real k1[6], k2[6], k3[6], k4[6];
264 yprev[1] = p->
phi[i];
266 yprev[3] = p->
ppar[i];
268 yprev[5] = p->
zeta[i];
276 B_dB[4] = p->
B_phi[i];
291 errflag =
mhd_eval(mhd_dmhd, yprev[0], yprev[1], yprev[2], t0,
295 step_gceom_mhd(k1, yprev, mass, charge, B_dB, E, mhd_dmhd,
301 for(
int j = 0; j < 6; j++) {
302 tempy[j] = yprev[j] + h[i]/2.0*k1[j];
307 t0 + h[i]/2.0, Bdata);
311 t0 + h[i]/2.0, Edata, Bdata);
314 errflag =
mhd_eval(mhd_dmhd, tempy[0], tempy[1], tempy[2],
319 step_gceom_mhd(k2, tempy, mass, charge, B_dB, E, mhd_dmhd,
322 for(
int j = 0; j < 6; j++) {
323 tempy[j] = yprev[j] + h[i]/2.0*k2[j];
328 t0 + h[i]/2.0, Bdata);
332 t0 + h[i]/2.0, Edata, Bdata);
335 errflag =
mhd_eval(mhd_dmhd, tempy[0], tempy[1], tempy[2],
340 step_gceom_mhd(k3, tempy, mass, charge, B_dB, E, mhd_dmhd,
343 for(
int j = 0; j < 6; j++) {
344 tempy[j] = yprev[j] + h[i]*k3[j];
353 t0 + h[i], Edata, Bdata);
356 errflag =
mhd_eval(mhd_dmhd, tempy[0], tempy[1], tempy[2],
361 step_gceom_mhd(k4, tempy, mass, charge, B_dB, E, mhd_dmhd,
364 for(
int j = 0; j < 6; j++) {
366 + h[i]/6.0 * (k1[j] + 2*k2[j] + 2*k3[j] + k4[j]);
370 if(!errflag && y[0] <= 0) {
373 else if(!errflag && fabs(y[4]) >=
CONST_C) {
376 else if(!errflag && y[4] < 0) {
414 p->
B_phi[i] = B_dB[4];
429 p->
theta[i] += atan2( (R0-axisrz[0]) * (p->
z[i]-axisrz[1])
430 - (z0-axisrz[1]) * (p->
r[i]-axisrz[0]),
431 (R0-axisrz[0]) * (p->
r[i]-axisrz[0])
432 + (z0-axisrz[1]) * (p->
z[i]-axisrz[1]) );
a5err mhd_eval(real mhd_dmhd[10], real r, real phi, real z, real t, int includemode, boozer_data *boozerdata, mhd_data *mhddata, B_field_data *Bdata)
Evaluate the needed quantities from MHD mode for orbit following.
Header file for particle.c.
void step_gc_rk4_mhd(particle_simd_gc *p, real *h, B_field_data *Bdata, E_field_data *Edata, boozer_data *boozer, mhd_data *mhd, int aldforce)
Integrate a guiding center step with RK4 with MHD modes present.
Struct representing NSIMD guiding center markers.