ASCOT5
Loading...
Searching...
No Matches
dist_rho6D.c
Go to the documentation of this file.
1
5#include <stdio.h>
6#include <stdlib.h>
7#include <math.h>
8#include "../ascot5.h"
9#include "../consts.h"
10#include "../physlib.h"
11#include "../gctransform.h"
12#include "dist_rho6D.h"
13
17size_t dist_rho6D_index(int i_rho, int i_theta, int i_phi, int i_pr, int i_pphi,
18 int i_pz, int i_time, int i_q, size_t step_7,
19 size_t step_6, size_t step_5, size_t step_4,
20 size_t step_3, size_t step_2, size_t step_1) {
21 return (size_t)(i_rho) * step_7
22 + (size_t)(i_theta) * step_6
23 + (size_t)(i_phi) * step_5
24 + (size_t)(i_pr) * step_4
25 + (size_t)(i_pphi) * step_3
26 + (size_t)(i_pz) * step_2
27 + (size_t)(i_time) * step_1
28 + (size_t)(i_q);
29}
30
35
36 size_t n_q = (size_t)(data->n_q);
37 size_t n_time = (size_t)(data->n_time);
38 size_t n_pz = (size_t)(data->n_pz);
39 size_t n_pphi = (size_t)(data->n_pphi);
40 size_t n_pr = (size_t)(data->n_pr);
41 size_t n_phi = (size_t)(data->n_phi);
42 size_t n_theta = (size_t)(data->n_theta);
43 data->step_7 = n_q * n_time * n_pz * n_pphi * n_pr * n_phi * n_theta;
44 data->step_6 = n_q * n_time * n_pz * n_pphi * n_pr * n_phi;
45 data->step_5 = n_q * n_time * n_pz * n_pphi * n_pr;
46 data->step_4 = n_q * n_time * n_pz * n_pphi;
47 data->step_3 = n_q * n_time * n_pz;
48 data->step_2 = n_q * n_time;
49 data->step_1 = n_q;
50
51 data->histogram = calloc(data->step_7 * (size_t)data->n_rho, sizeof(real));
52 return data->histogram == NULL;
53}
54
59 free(data->histogram);
60}
61
68 GPU_MAP_TO_DEVICE(
69 data->histogram[0:data->n_rho*data->n_theta*data->n_phi*data->n_pr*data->n_pphi*data->n_pz*data->n_time*data->n_q]
70 )
71}
72
85 particle_simd_fo* p_i) {
86
87#ifdef GPU
88 size_t index;
89 real weight;
90#else
91 size_t index[NSIMD];
92 real weight[NSIMD];
93#endif
94
95 GPU_PARALLEL_LOOP_ALL_LEVELS
96 for(int i = 0; i < p_f->n_mrk; i++) {
97 if(p_f->running[i]) {
98
99 int i_rho = floor((p_f->rho[i] - dist->min_rho)
100 / ((dist->max_rho - dist->min_rho)/dist->n_rho));
101
102 real phi = fmod(p_f->phi[i], 2*CONST_PI);
103 if(phi < 0) {
104 phi = phi + 2*CONST_PI;
105 }
106 int i_phi = floor((phi - dist->min_phi)
107 / ((dist->max_phi - dist->min_phi)/dist->n_phi));
108
109 real theta = fmod(p_f->theta[i], 2*CONST_PI);
110 if(theta < 0) {
111 theta += 2*CONST_PI;
112 }
113 int i_theta = floor((theta - dist->min_theta)
114 / ((dist->max_theta - dist->min_theta)
115 / dist->n_theta));
116
117 int i_pr = floor((p_f->p_r[i] - dist->min_pr)
118 / ((dist->max_pr - dist->min_pr) / dist->n_pr));
119
120 int i_pphi = floor((p_f->p_phi[i] - dist->min_pphi)
121 / ((dist->max_pphi - dist->min_pphi)
122 / dist->n_pphi));
123
124 int i_pz = floor((p_f->p_z[i] - dist->min_pz)
125 / ((dist->max_pz - dist->min_pz) / dist->n_pz));
126
127 int i_time = floor((p_f->time[i] - dist->min_time)
128 / ((dist->max_time - dist->min_time) / dist->n_time));
129
130 int i_q = floor((p_f->charge[i]/CONST_E - dist->min_q)
131 / ((dist->max_q - dist->min_q) / dist->n_q));
132
133 if(i_rho >= 0 && i_rho <= dist->n_rho - 1 &&
134 i_theta >=0 && i_theta <= dist->n_theta -1 &&
135 i_phi >=0 && i_phi <= dist->n_phi - 1 &&
136 i_pr >= 0 && i_pr <= dist->n_pr - 1 &&
137 i_pphi >= 0 && i_pphi <= dist->n_pphi - 1 &&
138 i_pz >= 0 && i_pz <= dist->n_pz - 1 &&
139 i_time >= 0 && i_time <= dist->n_time - 1 &&
140 i_q >= 0 && i_q <= dist->n_q - 1 ) {
141#ifdef GPU
142 index = dist_rho6D_index(
143 i_rho, i_theta, i_phi, i_pr, i_pphi, i_pz,
144 i_time, i_q, dist->step_7, dist->step_6, dist->step_5,
145 dist->step_4, dist->step_3, dist->step_2, dist->step_1);
146 weight = p_f->weight[i] * (p_f->time[i] - p_i->time[i]);
147 GPU_ATOMIC
148 dist->histogram[index] += weight;
149#else
150 index[i] = dist_rho6D_index(
151 i_rho, i_theta, i_phi, i_pr, i_pphi, i_pz,
152 i_time, i_q, dist->step_7, dist->step_6, dist->step_5,
153 dist->step_4, dist->step_3, dist->step_2, dist->step_1);
154 weight[i] = p_f->weight[i] * (p_f->time[i] - p_i->time[i]);
155#endif
156 }
157 }
158 }
159#ifndef GPU
160 for(int i = 0; i < p_f->n_mrk; i++) {
161 if(p_f->running[i] && index[i] >= 0 &&
162 index[i] < dist->step_7 * dist->n_rho) {
163 GPU_ATOMIC
164 dist->histogram[index[i]] += weight[i];
165 }
166 }
167#endif
168}
169
182 particle_simd_gc* p_i) {
183 real phi[NSIMD];
184 real theta[NSIMD];
185
186 int i_rho[NSIMD];
187 int i_theta[NSIMD];
188 int i_phi[NSIMD];
189 int i_pr[NSIMD];
190 int i_pphi[NSIMD];
191 int i_pz[NSIMD];
192 int i_time[NSIMD];
193 int i_q[NSIMD];
194
195 int ok[NSIMD];
196 real weight[NSIMD];
197
198 #pragma omp simd
199 for(int i = 0; i < NSIMD; i++) {
200 if(p_f->running[i]) {
201
202 real pr, pphi, pz;
203 real B_dB[12] = {p_f->B_r[i],
204 p_f->B_r_dr[i],
205 p_f->B_r_dphi[i],
206 p_f->B_r_dz[i],
207 p_f->B_phi[i],
208 p_f->B_phi_dr[i],
209 p_f->B_phi_dphi[i],
210 p_f->B_phi_dz[i],
211 p_f->B_z[i],
212 p_f->B_z_dr[i],
213 p_f->B_z_dphi[i],
214 p_f->B_z_dz[i]};
215 gctransform_pparmuzeta2prpphipz(p_f->mass[i], p_f->charge[i], B_dB,
216 p_f->phi[i], p_f->ppar[i],
217 p_f->mu[i], p_f->zeta[i],
218 &pr, &pphi, &pz);
219
220 i_rho[i] = floor((p_f->rho[i] - dist->min_rho)
221 / ((dist->max_rho - dist->min_rho)/dist->n_rho));
222
223 phi[i] = fmod(p_f->phi[i], 2*CONST_PI);
224 if(phi[i] < 0) {
225 phi[i] = phi[i] + 2*CONST_PI;
226 }
227 i_phi[i] = floor((phi[i] - dist->min_phi)
228 / ((dist->max_phi - dist->min_phi)/dist->n_phi));
229
230 theta[i] = fmod(p_f->theta[i], 2*CONST_PI);
231 if(theta[i] < 0) {
232 theta[i] = theta[i] + 2*CONST_PI;
233 }
234 i_theta[i] = floor((theta[i] - dist->min_theta)
235 / ((dist->max_theta - dist->min_theta)
236 / dist->n_theta));
237
238 i_pr[i] = floor((pr - dist->min_pr)
239 / ((dist->max_pr - dist->min_pr) / dist->n_pr));
240
241 i_pphi[i] = floor((pphi - dist->min_pphi)
242 / ((dist->max_pphi - dist->min_pphi)
243 / dist->n_pphi));
244
245 i_pz[i] = floor((pz - dist->min_pz)
246 / ((dist->max_pz - dist->min_pz) / dist->n_pz));
247
248 i_time[i] = floor((p_f->time[i] - dist->min_time)
249 / ((dist->max_time - dist->min_time) / dist->n_time));
250
251 i_q[i] = floor((p_f->charge[i]/CONST_E - dist->min_q)
252 / ((dist->max_q - dist->min_q) / dist->n_q));
253
254 if(i_rho[i] >= 0 && i_rho[i] <= dist->n_rho - 1 &&
255 i_theta[i] >= 0 && i_theta[i] <= dist->n_theta -1 &&
256 i_phi[i] >= 0 && i_phi[i] <= dist->n_phi - 1 &&
257 i_pr[i] >= 0 && i_pr[i] <= dist->n_pr - 1 &&
258 i_pphi[i] >= 0 && i_pphi[i] <= dist->n_pphi - 1 &&
259 i_pz[i] >= 0 && i_pz[i] <= dist->n_pz - 1 &&
260 i_time[i] >= 0 && i_time[i] <= dist->n_time - 1 &&
261 i_q[i] >= 0 && i_q[i] <= dist->n_q - 1 ) {
262 ok[i] = 1;
263 weight[i] = p_f->weight[i] * (p_f->time[i] - p_i->time[i]);
264 }
265 else {
266 ok[i] = 0;
267 }
268 }
269 }
270
271 for(int i = 0; i < NSIMD; i++) {
272 if(p_f->running[i] && ok[i]) {
273 size_t index = dist_rho6D_index(
274 i_rho[i], i_theta[i], i_phi[i], i_pr[i], i_pphi[i], i_pz[i],
275 i_time[i], i_q[i], dist->step_7, dist->step_6, dist->step_5,
276 dist->step_4, dist->step_3, dist->step_2, dist->step_1);
277 #pragma omp atomic
278 dist->histogram[index] += weight[i];
279 }
280 }
281}
Main header file for ASCOT5.
double real
Definition ascot5.h:85
#define NSIMD
Number of particles simulated simultaneously in a particle group operations.
Definition ascot5.h:91
Header file containing physical and mathematical constants.
#define CONST_PI
pi
Definition consts.h:11
#define CONST_E
Elementary charge [C].
Definition consts.h:35
int dist_rho6D_init(dist_rho6D_data *data)
Initializes distribution data.
Definition dist_rho6D.c:34
void dist_rho6D_update_gc(dist_rho6D_data *dist, particle_simd_gc *p_f, particle_simd_gc *p_i)
Update the histogram from guiding-center particles.
Definition dist_rho6D.c:181
size_t dist_rho6D_index(int i_rho, int i_theta, int i_phi, int i_pr, int i_pphi, int i_pz, int i_time, int i_q, size_t step_7, size_t step_6, size_t step_5, size_t step_4, size_t step_3, size_t step_2, size_t step_1)
Internal function calculating the index in the histogram array.
Definition dist_rho6D.c:17
void dist_rho6D_free(dist_rho6D_data *data)
Free allocated resources.
Definition dist_rho6D.c:58
void dist_rho6D_update_fo(dist_rho6D_data *dist, particle_simd_fo *p_f, particle_simd_fo *p_i)
Update the histogram from full-orbit particles.
Definition dist_rho6D.c:84
void dist_rho6D_offload(dist_rho6D_data *data)
Offload data to the accelerator.
Definition dist_rho6D.c:67
Header file for dist_rho6D.c.
void gctransform_pparmuzeta2prpphipz(real mass, real charge, real *B_dB, real phi, real ppar, real mu, real zeta, real *pr, real *pphi, real *pz)
Transform particle ppar, mu, and zeta to momentum vector.
Header file for gctransform.c.
real fmod(real x, real y)
Compute the modulus of two real numbers.
Definition math.c:22
Header file for math.c.
Methods to evaluate elementary physical quantities.
Histogram parameters on target.
Definition dist_rho6D.h:15
Struct representing NSIMD particle markers.
Definition particle.h:210
integer * running
Definition particle.h:252
Struct representing NSIMD guiding center markers.
Definition particle.h:275