MCBooster  1.0.1
Tool to generate MC phase space samples in parallel.
EvaluateArray.h
Go to the documentation of this file.
1 /*
2  * EvaluateArray.h
3  *
4  * Copyright 2016 Antonio Augusto Alves Junior
5  *
6  * Created on : Feb 25, 2016
7  * Author: Antonio Augusto Alves Junior
8  */
9 
10 /*
11  * This file is part of MCBooster.
12  *
13  * MCBooster is free software: you can redistribute it and/or modify
14  * it under the terms of the GNU General Public License as published by
15  * the Free Software Foundation, either version 3 of the License, or
16  * (at your option) any later version.
17  *
18  * MCBooster is distributed in the hope that it will be useful,
19  * but WITHOUT ANY WARRANTY; without even the implied warranty of
20  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
21  * GNU General Public License for more details.
22  *
23  * You should have received a copy of the GNU General Public License
24  * along with MCBooster. If not, see <http://www.gnu.org/licenses/>.
25  */
30 #ifndef EVALUATEARRAY_H_
31 #define EVALUATEARRAY_H_
32 
33 
34 #include <mcbooster/Config.h>
35 #include <mcbooster/Vector3R.h>
36 #include <mcbooster/Vector4R.h>
37 #include <mcbooster/GContainers.h>
38 #include <mcbooster/GTypes.h>
39 #include <mcbooster/functors/Calculate.h>
40 
41 namespace MCBooster
42 {
43 
50 template<typename CUSTOMFUNC>
51 void EvaluateArray(const CUSTOMFUNC funcObj, ParticlesSet_d &pset,
52  VariableSet_h &varset)
53 {
54 
55  if (pset.size() > 10 || pset.size() < 2)
56  {
57  cout
58  << "Can not Calculate(Eval) more than a nine-particle invariant mass."
59  << endl;
60  return;
61  }
62 
63  GInt_t arrayWidth = varset.size();
64  GLong_t numberEvents = varset[0]->size();
65 
66  RealVector_d dev_array(arrayWidth * numberEvents);
67 
68  strided_range<RealVector_d::iterator> it_array(dev_array.begin(),
69  dev_array.end(), arrayWidth);
70 
71  if (pset.size() == 2)
72  {
73 
74  mc_device_vector<Vector4R> *dev_v0 = pset[0];
75  mc_device_vector<Vector4R> *dev_v1 = pset[1];
76 
77  thrust::for_each(
78  thrust::make_zip_iterator(
79  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
80  dev_v1->begin())),
81  thrust::make_zip_iterator(
82  thrust::make_tuple(it_array.end(), dev_v0->end(),
83  dev_v1->end())),
84  Calculate3<CUSTOMFUNC>(funcObj));
85 
86  }
87 
88  if (pset.size() == 3)
89  {
90 
91  mc_device_vector<Vector4R> *dev_v0 = pset[0];
92  mc_device_vector<Vector4R> *dev_v1 = pset[1];
93  mc_device_vector<Vector4R> *dev_v2 = pset[2];
94 
95  thrust::for_each(
96  thrust::make_zip_iterator(
97  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
98  dev_v1->begin(), dev_v2->begin())),
99  thrust::make_zip_iterator(
100  thrust::make_tuple(it_array.end(), dev_v0->end(),
101  dev_v1->end(), dev_v2->end())),
102  Calculate3<CUSTOMFUNC>(funcObj));
103 
104  }
105 
106  if (pset.size() == 4)
107  {
108 
109  mc_device_vector<Vector4R> *dev_v0 = pset[0];
110  mc_device_vector<Vector4R> *dev_v1 = pset[1];
111  mc_device_vector<Vector4R> *dev_v2 = pset[2];
112  mc_device_vector<Vector4R> *dev_v3 = pset[3];
113 
114  thrust::for_each(
115  thrust::make_zip_iterator(
116  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
117  dev_v1->begin(), dev_v2->begin(),
118  dev_v3->begin())),
119  thrust::make_zip_iterator(
120  thrust::make_tuple(it_array.end(), dev_v0->end(),
121  dev_v1->end(), dev_v2->end(), dev_v3->end())),
122  Calculate3<CUSTOMFUNC>(funcObj));
123 
124  }
125 
126  if (pset.size() == 5)
127  {
128 
129  mc_device_vector<Vector4R> *dev_v0 = pset[0];
130  mc_device_vector<Vector4R> *dev_v1 = pset[1];
131  mc_device_vector<Vector4R> *dev_v2 = pset[2];
132  mc_device_vector<Vector4R> *dev_v3 = pset[3];
133  mc_device_vector<Vector4R> *dev_v4 = pset[4];
134 
135  thrust::for_each(
136  thrust::make_zip_iterator(
137  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
138  dev_v1->begin(), dev_v2->begin(),
139  dev_v3->begin(), dev_v4->begin())),
140  thrust::make_zip_iterator(
141  thrust::make_tuple(it_array.end(), dev_v0->end(),
142  dev_v1->end(), dev_v2->end(), dev_v3->end(),
143  dev_v4->end())),
144  Calculate3<CUSTOMFUNC>(funcObj));
145 
146  }
147 
148  if (pset.size() == 6)
149  {
150 
151  mc_device_vector<Vector4R> *dev_v0 = pset[0];
152  mc_device_vector<Vector4R> *dev_v1 = pset[1];
153  mc_device_vector<Vector4R> *dev_v2 = pset[2];
154  mc_device_vector<Vector4R> *dev_v3 = pset[3];
155  mc_device_vector<Vector4R> *dev_v4 = pset[4];
156  mc_device_vector<Vector4R> *dev_v5 = pset[5];
157 
158  thrust::for_each(
159  thrust::make_zip_iterator(
160  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
161  dev_v1->begin(), dev_v2->begin(),
162  dev_v3->begin(), dev_v4->begin(),
163  dev_v5->begin())),
164  thrust::make_zip_iterator(
165  thrust::make_tuple(it_array.end(), dev_v0->end(),
166  dev_v1->end(), dev_v2->end(), dev_v3->end(),
167  dev_v4->end(), dev_v5->end())),
168  Calculate3<CUSTOMFUNC>(funcObj));
169 
170  }
171 
172  if (pset.size() == 7)
173  {
174 
175  mc_device_vector<Vector4R> *dev_v0 = pset[0];
176  mc_device_vector<Vector4R> *dev_v1 = pset[1];
177  mc_device_vector<Vector4R> *dev_v2 = pset[2];
178  mc_device_vector<Vector4R> *dev_v3 = pset[3];
179  mc_device_vector<Vector4R> *dev_v4 = pset[4];
180  mc_device_vector<Vector4R> *dev_v5 = pset[5];
181  mc_device_vector<Vector4R> *dev_v6 = pset[6];
182 
183  thrust::for_each(
184  thrust::make_zip_iterator(
185  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
186  dev_v1->begin(), dev_v2->begin(),
187  dev_v3->begin(), dev_v4->begin(),
188  dev_v5->begin(), dev_v6->begin())),
189  thrust::make_zip_iterator(
190  thrust::make_tuple(it_array.end(), dev_v0->end(),
191  dev_v1->end(), dev_v2->end(), dev_v3->end(),
192  dev_v4->end(), dev_v5->end(), dev_v6->end())),
193  Calculate3<CUSTOMFUNC>(funcObj));
194 
195  }
196 
197  if (pset.size() == 8)
198  {
199 
200  mc_device_vector<Vector4R> *dev_v0 = pset[0];
201  mc_device_vector<Vector4R> *dev_v1 = pset[1];
202  mc_device_vector<Vector4R> *dev_v2 = pset[2];
203  mc_device_vector<Vector4R> *dev_v3 = pset[3];
204  mc_device_vector<Vector4R> *dev_v4 = pset[4];
205  mc_device_vector<Vector4R> *dev_v5 = pset[5];
206  mc_device_vector<Vector4R> *dev_v6 = pset[6];
207  mc_device_vector<Vector4R> *dev_v7 = pset[7];
208 
209  thrust::for_each(
210  thrust::make_zip_iterator(
211  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
212  dev_v1->begin(), dev_v2->begin(),
213  dev_v3->begin(), dev_v4->begin(),
214  dev_v5->begin(), dev_v6->begin(),
215  dev_v7->begin())),
216  thrust::make_zip_iterator(
217  thrust::make_tuple(it_array.end(), dev_v0->end(),
218  dev_v1->end(), dev_v2->end(), dev_v3->end(),
219  dev_v4->end(), dev_v5->end(), dev_v6->end(),
220  dev_v7->end())),
221  Calculate3<CUSTOMFUNC>(funcObj));
222 
223  }
224 
225  if (pset.size() == 9)
226  {
227 
228  mc_device_vector<Vector4R> *dev_v0 = pset[0];
229  mc_device_vector<Vector4R> *dev_v1 = pset[1];
230  mc_device_vector<Vector4R> *dev_v2 = pset[2];
231  mc_device_vector<Vector4R> *dev_v3 = pset[3];
232  mc_device_vector<Vector4R> *dev_v4 = pset[4];
233  mc_device_vector<Vector4R> *dev_v5 = pset[5];
234  mc_device_vector<Vector4R> *dev_v6 = pset[6];
235  mc_device_vector<Vector4R> *dev_v7 = pset[7];
236  mc_device_vector<Vector4R> *dev_v8 = pset[8];
237 
238  thrust::for_each(
239  thrust::make_zip_iterator(
240  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
241  dev_v1->begin(), dev_v2->begin(),
242  dev_v3->begin(), dev_v4->begin(),
243  dev_v5->begin(), dev_v6->begin(),
244  dev_v7->begin(), dev_v8->begin())),
245  thrust::make_zip_iterator(
246  thrust::make_tuple(it_array.end(), dev_v0->end(),
247  dev_v1->end(), dev_v2->end(), dev_v3->end(),
248  dev_v4->end(), dev_v5->end(), dev_v6->end(),
249  dev_v7->end(), dev_v8->end())),
250  Calculate3<CUSTOMFUNC>(funcObj));
251 
252  }
253 
254  if (pset.size() == 10)
255  {
256 
257  mc_device_vector<Vector4R> *dev_v0 = pset[0];
258  mc_device_vector<Vector4R> *dev_v1 = pset[1];
259  mc_device_vector<Vector4R> *dev_v2 = pset[2];
260  mc_device_vector<Vector4R> *dev_v3 = pset[3];
261  mc_device_vector<Vector4R> *dev_v4 = pset[4];
262  mc_device_vector<Vector4R> *dev_v5 = pset[5];
263  mc_device_vector<Vector4R> *dev_v6 = pset[6];
264  mc_device_vector<Vector4R> *dev_v7 = pset[7];
265  mc_device_vector<Vector4R> *dev_v8 = pset[8];
266 
267  thrust::for_each(
268  thrust::make_zip_iterator(
269  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
270  dev_v1->begin(), dev_v2->begin(),
271  dev_v3->begin(), dev_v4->begin(),
272  dev_v5->begin(), dev_v6->begin(),
273  dev_v7->begin(), dev_v8->begin())),
274  thrust::make_zip_iterator(
275  thrust::make_tuple(it_array.end(), dev_v0->end(),
276  dev_v1->end(), dev_v2->end(), dev_v3->end(),
277  dev_v4->end(), dev_v5->end(), dev_v6->end(),
278  dev_v7->end(), dev_v8->end())),
279  Calculate3<CUSTOMFUNC>(funcObj));
280 
281  }
282 
283 #if THRUST_DEVICE_SYSTEM==THRUST_DEVICE_BACKEND_OMP || THRUST_DEVICE_SYSTEM==THRUST_DEVICE_BACKEND_TBB
284 
285 #pragma omp parallel num_threads( arrayWidth )
286  {
287  strided_range<RealVector_d::iterator> it_array(dev_array.begin() + omp_get_thread_num()
288  , dev_array.end(), arrayWidth);
289 
290  thrust::copy(it_array.begin(),it_array.end(),
291  varset[omp_get_thread_num()]->begin());
292  }
293 #else
294  cudaStream_t s[arrayWidth];
295 
296  for (GInt_t d = 0; d < arrayWidth; d++)
297  {
298  cudaStreamCreate(&s[d]);
299 
300  }
301  strided_range<RealVector_d::iterator> *it[arrayWidth];
302  for (GInt_t d = 0; d < arrayWidth; d++)
303  it[d] = new strided_range<RealVector_d::iterator>(dev_array.begin() + d,
304  dev_array.end(), arrayWidth);
305  for (GInt_t d = 0; d < arrayWidth; d++)
306  {
307 
308  thrust::copy(thrust::cuda::par.on(s[d]), it[d]->begin(), it[d]->end(),
309  varset[d]->begin());
310 
311  }
312  cudaDeviceSynchronize();
313  for (GInt_t d = 0; d < arrayWidth; d++)
314  cudaStreamDestroy(s[d]);
315  for (GInt_t d = 0; d < arrayWidth; d++)
316  delete it[d];
317 
318 #endif
319  return;
320 }
321 
322 #if !(THRUST_DEVICE_SYSTEM==THRUST_DEVICE_BACKEND_OMP || THRUST_DEVICE_SYSTEM==THRUST_DEVICE_BACKEND_TBB)
323 
331 template<typename CUSTOMFUNC>
332 void EvaluateArray(const CUSTOMFUNC funcObj, ParticlesSet_d &pset,
333  VariableSet_d &varset)
334 {
335 
336  if (pset.size() > 10 || pset.size() < 2)
337  {
338  cout
339  << "Can not Calculate(Eval) more than a nine-particle invariant mass."
340  << endl;
341  return;
342  }
343 
344  GInt_t arrayWidth = varset.size();
345  GLong_t numberEvents = varset[0]->size();
346 
347  RealVector_d dev_array(arrayWidth * numberEvents);
348 
349  strided_range<RealVector_d::iterator> it_array(dev_array.begin(),
350  dev_array.end(), arrayWidth);
351 
352  if (pset.size() == 2)
353  {
354 
355  mc_device_vector<Vector4R> *dev_v0 = pset[0];
356  mc_device_vector<Vector4R> *dev_v1 = pset[1];
357 
358  thrust::for_each(
359  thrust::make_zip_iterator(
360  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
361  dev_v1->begin())),
362  thrust::make_zip_iterator(
363  thrust::make_tuple(it_array.end(), dev_v0->end(),
364  dev_v1->end())),
365  Calculate3<CUSTOMFUNC>(funcObj));
366 
367  }
368 
369  if (pset.size() == 3)
370  {
371 
372  mc_device_vector<Vector4R> *dev_v0 = pset[0];
373  mc_device_vector<Vector4R> *dev_v1 = pset[1];
374  mc_device_vector<Vector4R> *dev_v2 = pset[2];
375 
376  thrust::for_each(
377  thrust::make_zip_iterator(
378  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
379  dev_v1->begin(), dev_v2->begin())),
380  thrust::make_zip_iterator(
381  thrust::make_tuple(it_array.end(), dev_v0->end(),
382  dev_v1->end(), dev_v2->end())),
383  Calculate3<CUSTOMFUNC>(funcObj));
384 
385  }
386 
387  if (pset.size() == 4)
388  {
389 
390  mc_device_vector<Vector4R> *dev_v0 = pset[0];
391  mc_device_vector<Vector4R> *dev_v1 = pset[1];
392  mc_device_vector<Vector4R> *dev_v2 = pset[2];
393  mc_device_vector<Vector4R> *dev_v3 = pset[3];
394 
395  thrust::for_each(
396  thrust::make_zip_iterator(
397  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
398  dev_v1->begin(), dev_v2->begin(),
399  dev_v3->begin())),
400  thrust::make_zip_iterator(
401  thrust::make_tuple(it_array.end(), dev_v0->end(),
402  dev_v1->end(), dev_v2->end(), dev_v3->end())),
403  Calculate3<CUSTOMFUNC>(funcObj));
404 
405  }
406 
407  if (pset.size() == 5)
408  {
409 
410  mc_device_vector<Vector4R> *dev_v0 = pset[0];
411  mc_device_vector<Vector4R> *dev_v1 = pset[1];
412  mc_device_vector<Vector4R> *dev_v2 = pset[2];
413  mc_device_vector<Vector4R> *dev_v3 = pset[3];
414  mc_device_vector<Vector4R> *dev_v4 = pset[4];
415 
416  thrust::for_each(
417  thrust::make_zip_iterator(
418  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
419  dev_v1->begin(), dev_v2->begin(),
420  dev_v3->begin(), dev_v4->begin())),
421  thrust::make_zip_iterator(
422  thrust::make_tuple(it_array.end(), dev_v0->end(),
423  dev_v1->end(), dev_v2->end(), dev_v3->end(),
424  dev_v4->end())),
425  Calculate3<CUSTOMFUNC>(funcObj));
426 
427  }
428 
429  if (pset.size() == 6)
430  {
431 
432  mc_device_vector<Vector4R> *dev_v0 = pset[0];
433  mc_device_vector<Vector4R> *dev_v1 = pset[1];
434  mc_device_vector<Vector4R> *dev_v2 = pset[2];
435  mc_device_vector<Vector4R> *dev_v3 = pset[3];
436  mc_device_vector<Vector4R> *dev_v4 = pset[4];
437  mc_device_vector<Vector4R> *dev_v5 = pset[5];
438 
439  thrust::for_each(
440  thrust::make_zip_iterator(
441  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
442  dev_v1->begin(), dev_v2->begin(),
443  dev_v3->begin(), dev_v4->begin(),
444  dev_v5->begin())),
445  thrust::make_zip_iterator(
446  thrust::make_tuple(it_array.end(), dev_v0->end(),
447  dev_v1->end(), dev_v2->end(), dev_v3->end(),
448  dev_v4->end(), dev_v5->end())),
449  Calculate3<CUSTOMFUNC>(funcObj));
450 
451  }
452 
453  if (pset.size() == 7)
454  {
455 
456  mc_device_vector<Vector4R> *dev_v0 = pset[0];
457  mc_device_vector<Vector4R> *dev_v1 = pset[1];
458  mc_device_vector<Vector4R> *dev_v2 = pset[2];
459  mc_device_vector<Vector4R> *dev_v3 = pset[3];
460  mc_device_vector<Vector4R> *dev_v4 = pset[4];
461  mc_device_vector<Vector4R> *dev_v5 = pset[5];
462  mc_device_vector<Vector4R> *dev_v6 = pset[6];
463 
464  thrust::for_each(
465  thrust::make_zip_iterator(
466  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
467  dev_v1->begin(), dev_v2->begin(),
468  dev_v3->begin(), dev_v4->begin(),
469  dev_v5->begin(), dev_v6->begin())),
470  thrust::make_zip_iterator(
471  thrust::make_tuple(it_array.end(), dev_v0->end(),
472  dev_v1->end(), dev_v2->end(), dev_v3->end(),
473  dev_v4->end(), dev_v5->end(), dev_v6->end())),
474  Calculate3<CUSTOMFUNC>(funcObj));
475 
476  }
477 
478  if (pset.size() == 8)
479  {
480 
481  mc_device_vector<Vector4R> *dev_v0 = pset[0];
482  mc_device_vector<Vector4R> *dev_v1 = pset[1];
483  mc_device_vector<Vector4R> *dev_v2 = pset[2];
484  mc_device_vector<Vector4R> *dev_v3 = pset[3];
485  mc_device_vector<Vector4R> *dev_v4 = pset[4];
486  mc_device_vector<Vector4R> *dev_v5 = pset[5];
487  mc_device_vector<Vector4R> *dev_v6 = pset[6];
488  mc_device_vector<Vector4R> *dev_v7 = pset[7];
489 
490  thrust::for_each(
491  thrust::make_zip_iterator(
492  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
493  dev_v1->begin(), dev_v2->begin(),
494  dev_v3->begin(), dev_v4->begin(),
495  dev_v5->begin(), dev_v6->begin(),
496  dev_v7->begin())),
497  thrust::make_zip_iterator(
498  thrust::make_tuple(it_array.end(), dev_v0->end(),
499  dev_v1->end(), dev_v2->end(), dev_v3->end(),
500  dev_v4->end(), dev_v5->end(), dev_v6->end(),
501  dev_v7->end())),
502  Calculate3<CUSTOMFUNC>(funcObj));
503 
504  }
505 
506  if (pset.size() == 9)
507  {
508 
509  mc_device_vector<Vector4R> *dev_v0 = pset[0];
510  mc_device_vector<Vector4R> *dev_v1 = pset[1];
511  mc_device_vector<Vector4R> *dev_v2 = pset[2];
512  mc_device_vector<Vector4R> *dev_v3 = pset[3];
513  mc_device_vector<Vector4R> *dev_v4 = pset[4];
514  mc_device_vector<Vector4R> *dev_v5 = pset[5];
515  mc_device_vector<Vector4R> *dev_v6 = pset[6];
516  mc_device_vector<Vector4R> *dev_v7 = pset[7];
517  mc_device_vector<Vector4R> *dev_v8 = pset[8];
518 
519  thrust::for_each(
520  thrust::make_zip_iterator(
521  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
522  dev_v1->begin(), dev_v2->begin(),
523  dev_v3->begin(), dev_v4->begin(),
524  dev_v5->begin(), dev_v6->begin(),
525  dev_v7->begin(), dev_v8->begin())),
526  thrust::make_zip_iterator(
527  thrust::make_tuple(it_array.end(), dev_v0->end(),
528  dev_v1->end(), dev_v2->end(), dev_v3->end(),
529  dev_v4->end(), dev_v5->end(), dev_v6->end(),
530  dev_v7->end(), dev_v8->end())),
531  Calculate3<CUSTOMFUNC>(funcObj));
532 
533  }
534 
535  if (pset.size() == 10)
536  {
537 
538  mc_device_vector<Vector4R> *dev_v0 = pset[0];
539  mc_device_vector<Vector4R> *dev_v1 = pset[1];
540  mc_device_vector<Vector4R> *dev_v2 = pset[2];
541  mc_device_vector<Vector4R> *dev_v3 = pset[3];
542  mc_device_vector<Vector4R> *dev_v4 = pset[4];
543  mc_device_vector<Vector4R> *dev_v5 = pset[5];
544  mc_device_vector<Vector4R> *dev_v6 = pset[6];
545  mc_device_vector<Vector4R> *dev_v7 = pset[7];
546  mc_device_vector<Vector4R> *dev_v8 = pset[8];
547 
548  thrust::for_each(
549  thrust::make_zip_iterator(
550  thrust::make_tuple(it_array.begin(), dev_v0->begin(),
551  dev_v1->begin(), dev_v2->begin(),
552  dev_v3->begin(), dev_v4->begin(),
553  dev_v5->begin(), dev_v6->begin(),
554  dev_v7->begin(), dev_v8->begin())),
555  thrust::make_zip_iterator(
556  thrust::make_tuple(it_array.end(), dev_v0->end(),
557  dev_v1->end(), dev_v2->end(), dev_v3->end(),
558  dev_v4->end(), dev_v5->end(), dev_v6->end(),
559  dev_v7->end(), dev_v8->end())),
560  Calculate3<CUSTOMFUNC>(funcObj));
561 
562  }
563 
564 #if THRUST_DEVICE_SYSTEM==THRUST_DEVICE_BACKEND_OMP || THRUST_DEVICE_SYSTEM==THRUST_DEVICE_BACKEND_TBB
565 
566 #pragma omp parallel num_threads( arrayWidth )
567  {
568  strided_range<RealVector_d::iterator> it_array(dev_array.begin() + omp_get_thread_num()
569  , dev_array.end(), arrayWidth);
570 
571  thrust::copy(it_array.begin(),it_array.end(),
572  varset[omp_get_thread_num()]->begin());
573  }
574 #else
575  cudaStream_t s[arrayWidth];
576 
577  for (GInt_t d = 0; d < arrayWidth; d++)
578  {
579  cudaStreamCreate(&s[d]);
580 
581  }
582  strided_range<RealVector_d::iterator> *it[arrayWidth];
583  for (GInt_t d = 0; d < arrayWidth; d++)
584  it[d] = new strided_range<RealVector_d::iterator>(dev_array.begin() + d,
585  dev_array.end(), arrayWidth);
586  for (GInt_t d = 0; d < arrayWidth; d++)
587  {
588 
589  thrust::copy(thrust::cuda::par.on(s[d]), it[d]->begin(), it[d]->end(),
590  varset[d]->begin());
591 
592  }
593  cudaDeviceSynchronize();
594  for (GInt_t d = 0; d < arrayWidth; d++)
595  cudaStreamDestroy(s[d]);
596  for (GInt_t d = 0; d < arrayWidth; d++)
597  delete it[d];
598 
599 #endif
600  return;
601 }
602 #endif
603 
604 }
605 
606 #endif /* EVALUATEARRAY_H_ */
Vector4R.h
MCBooster::GLong_t
long GLong_t
Signed long integer 4 bytes (long)
Definition: GTypes.h:39
MCBooster::EvaluateArray
void EvaluateArray(const CUSTOMFUNC funcObj, ParticlesSet_d &pset, VariableSet_h &varset)
Template functor for calculate an array of variables over a given set of particles.
Definition: EvaluateArray.h:51
strided_range
Strided range iterator original code: https://github.com/thrust/thrust/blob/master/examples/strided_r...
Definition: strided_iterator.h:40
GContainers.h
Typedef for useful container classes used in MCBooster.
Calculate.h
MCBooster::VariableSet_d
vector< RealVector_d * > VariableSet_d
Typedef for a STL vector of pointers to device Particles_d vectors.
Definition: GContainers.h:117
Vector3R.h
GTypes.h
Config.h
MCBooster::ParticlesSet_d
vector< Particles_d * > ParticlesSet_d
Typedef for a Vector4R device vector.
Definition: GContainers.h:116
MCBooster::VariableSet_h
vector< RealVector_h * > VariableSet_h
Typedef for a STL vector of pointers to host Particles_h vectors .
Definition: GContainers.h:107
MCBooster::GInt_t
int GInt_t
Signed integer 4 bytes (int)
Definition: GTypes.h:37
MCBooster::mc_device_vector
thrust::host_vector< T > mc_device_vector
Generic template typedef for thrust::host_vector.
Definition: GContainers.h:60
MCBooster::RealVector_d
mc_device_vector< GReal_t > RealVector_d
Typedef for a GBool_t device vector.
Definition: GContainers.h:113