SpatialOps
NeboApplyPointwise.h
1 /* This file was generated by fulmar version 0.9.2. */
2 
3 #ifndef FULMAR_APPLY_POINTWISE_H
4  #define FULMAR_APPLY_POINTWISE_H
5 
6  namespace SpatialOps {
7  template<typename CurrentMode, typename FunctorFactory, typename Operand1>
9  template<typename FunctorFactory, typename Operand1>
10  struct ApplyPointwise1<Initial, FunctorFactory, Operand1> {
11  public:
12  ApplyPointwise1<SeqWalk,
13  FunctorFactory,
14  typename Operand1::SeqWalkType> typedef SeqWalkType;
15 
16  #ifdef ENABLE_THREADS
17  ApplyPointwise1<Resize,
18  FunctorFactory,
19  typename Operand1::ResizeType> typedef ResizeType;
20  #endif
21  /* ENABLE_THREADS */
22 
23  #ifdef __CUDACC__
24  ApplyPointwise1<GPUWalk,
25  FunctorFactory,
26  typename Operand1::GPUWalkType> typedef GPUWalkType
27  ;
28  #endif
29  /* __CUDACC__ */
30 
31  ApplyPointwise1(FunctorFactory const & functorFactory,
32  Operand1 const & operand1)
33  : functorFactory_(functorFactory), operand1_(operand1)
34  {}
35 
36  inline GhostData ghosts_with_bc(void) const {
37  return operand1_.ghosts_with_bc();
38  }
39 
40  inline GhostData ghosts_without_bc(void) const {
41  return operand1_.ghosts_without_bc();
42  }
43 
44  inline bool has_extents(void) const {
45  return (operand1_.has_extents());
46  }
47 
48  inline IntVec extents(void) const { return operand1_.extents(); }
49 
50  inline IntVec has_bc(void) const { return operand1_.has_bc(); }
51 
52  inline SeqWalkType init(IntVec const & extents,
53  GhostData const & ghosts,
54  IntVec const & hasBC) const {
55  return SeqWalkType(functorFactory_,
56  operand1_.init(extents, ghosts, hasBC));
57  }
58 
59  #ifdef ENABLE_THREADS
60  inline ResizeType resize(void) const {
61  return ResizeType(functorFactory_, operand1_.resize());
62  }
63  #endif
64  /* ENABLE_THREADS */
65 
66  #ifdef __CUDACC__
67  inline bool cpu_ready(void) const {
68  return (operand1_.cpu_ready());
69  }
70 
71  inline bool gpu_ready(int const deviceIndex) const { return false; }
72 
73  inline GPUWalkType gpu_init(IntVec const & extents,
74  GhostData const & ghosts,
75  IntVec const & hasBC,
76  int const deviceIndex,
77  cudaStream_t const & lhsStream) const {
78  return GPUWalkType(operand1_.gpu_init(extents,
79  ghosts,
80  hasBC,
81  deviceIndex,
82  lhsStream));
83  }
84 
85  inline void stream_wait_event(cudaEvent_t const & event) const {
86  operand1_.stream_wait_event(event);
87  }
88 
89  #ifdef NEBO_GPU_TEST
90  inline void gpu_prep(int const deviceIndex) const {
91  operand1_.gpu_prep(deviceIndex);
92  }
93  #endif
94  /* NEBO_GPU_TEST */
95  #endif
96  /* __CUDACC__ */
97 
98  private:
99  FunctorFactory const & functorFactory_;
100 
101  Operand1 const operand1_;
102  };
103  #ifdef ENABLE_THREADS
104  template<typename FunctorFactory, typename Operand1>
105  struct ApplyPointwise1<Resize, FunctorFactory, Operand1> {
106  public:
107  ApplyPointwise1<SeqWalk,
108  FunctorFactory,
109  typename Operand1::SeqWalkType> typedef SeqWalkType
110  ;
111 
112  ApplyPointwise1(FunctorFactory const & functorFactory,
113  Operand1 const & operand1)
114  : functorFactory_(functorFactory), operand1_(operand1)
115  {}
116 
117  inline SeqWalkType init(IntVec const & extents,
118  GhostData const & ghosts,
119  IntVec const & hasBC) const {
120  return SeqWalkType(functorFactory_,
121  operand1_.init(extents, ghosts, hasBC));
122  }
123 
124  private:
125  FunctorFactory const & functorFactory_;
126 
127  Operand1 const operand1_;
128  }
129  #endif
130  /* ENABLE_THREADS */;
131  template<typename FunctorFactory, typename Operand1>
132  struct ApplyPointwise1<SeqWalk, FunctorFactory, Operand1> {
133  public:
134  typename Operand1::value_type typedef value_type;
135 
136  ApplyPointwise1(FunctorFactory const & functorFactory,
137  Operand1 const & operand1)
138  : functor_(functorFactory()), operand1_(operand1)
139  {}
140 
141  inline value_type eval(int const x, int const y, int const z) const {
142  return functor_(operand1_.eval(x, y, z));
143  }
144 
145  private:
146  typename FunctorFactory::Operator functor_;
147 
148  Operand1 operand1_;
149  };
150  #ifdef __CUDACC__
151  template<typename FunctorFactory, typename Operand1>
152  struct ApplyPointwise1<GPUWalk, FunctorFactory, Operand1> {
153  public:
154  typename Operand1::value_type typedef value_type;
155 
156  ApplyPointwise1(Operand1 const & operand1)
157  : operand1_(operand1)
158  {}
159 
160  __device__ inline value_type eval(int const x,
161  int const y,
162  int const z) const { return 0; }
163 
164  private:
165  Operand1 operand1_;
166  }
167  #endif
168  /* __CUDACC__ */;
169 
170  template<typename Operand1>
171  struct FindFieldType1 {
172  typename FindFieldType<Operand1>::Result typedef Result;
173  };
174 
175  template<typename FunctorFactory, typename Operand1>
176  inline NeboExpression<ApplyPointwise1<Initial,
177  FunctorFactory,
178  typename FinalType<Operand1>::
179  Result>,
181  apply_pointwise(FunctorFactory const & functorFactory, Operand1 operand1) {
182  ApplyPointwise1<Initial,
183  FunctorFactory,
184  typename FinalType<Operand1>::Result> typedef
185  ReturnType;
186 
188  typedef ReturnTerm;
189 
190  return ReturnTerm(ReturnType(functorFactory, normalize(operand1)));
191  }
192 
193  template<typename CurrentMode,
194  typename FunctorFactory,
195  typename Operand1,
196  typename Operand2>
198  template<typename FunctorFactory, typename Operand1, typename Operand2>
199  struct ApplyPointwise2<Initial, FunctorFactory, Operand1, Operand2> {
200  public:
201  ApplyPointwise2<SeqWalk,
202  FunctorFactory,
203  typename Operand1::SeqWalkType,
204  typename Operand2::SeqWalkType> typedef SeqWalkType;
205 
206  #ifdef ENABLE_THREADS
207  ApplyPointwise2<Resize,
208  FunctorFactory,
209  typename Operand1::ResizeType,
210  typename Operand2::ResizeType> typedef ResizeType;
211  #endif
212  /* ENABLE_THREADS */
213 
214  #ifdef __CUDACC__
215  ApplyPointwise2<GPUWalk,
216  FunctorFactory,
217  typename Operand1::GPUWalkType,
218  typename Operand2::GPUWalkType> typedef GPUWalkType
219  ;
220  #endif
221  /* __CUDACC__ */
222 
223  ApplyPointwise2(FunctorFactory const & functorFactory,
224  Operand1 const & operand1,
225  Operand2 const & operand2)
226  : functorFactory_(functorFactory), operand1_(operand1), operand2_(operand2)
227  {}
228 
229  inline GhostData ghosts_with_bc(void) const {
230  return min(operand1_.ghosts_with_bc(), operand2_.ghosts_with_bc());
231  }
232 
233  inline GhostData ghosts_without_bc(void) const {
234  return min(operand1_.ghosts_without_bc(), operand2_.ghosts_without_bc());
235  }
236 
237  inline bool has_extents(void) const {
238  return (operand1_.has_extents() || operand2_.has_extents());
239  }
240 
241  inline IntVec extents(void) const {
242  #ifndef NDEBUG
243  if((operand1_.has_extents() || operand2_.has_extents())) {
244  IntVec extents;
245 
246  if(operand1_.has_extents()) { extents = operand1_.extents(); }
247  else { extents = operand2_.extents(); };
248 
249  if(operand1_.has_extents()) {
250  assert(extents == operand1_.extents());
251  };
252 
253  if(operand2_.has_extents()) {
254  assert(extents == operand2_.extents());
255  };
256  }
257  #endif
258  /* NDEBUG */;
259 
260  return (operand1_.has_extents() ? operand1_.extents() : (operand2_.extents()));
261  }
262 
263  inline IntVec has_bc(void) const {
264  return (operand1_.has_extents() ? operand1_.has_bc() : (operand2_.has_bc()));
265  }
266 
267  inline SeqWalkType init(IntVec const & extents,
268  GhostData const & ghosts,
269  IntVec const & hasBC) const {
270  return SeqWalkType(functorFactory_,
271  operand1_.init(extents, ghosts, hasBC),
272  operand2_.init(extents, ghosts, hasBC));
273  }
274 
275  #ifdef ENABLE_THREADS
276  inline ResizeType resize(void) const {
277  return ResizeType(functorFactory_, operand1_.resize(), operand2_.resize());
278  }
279  #endif
280  /* ENABLE_THREADS */
281 
282  #ifdef __CUDACC__
283  inline bool cpu_ready(void) const {
284  return (operand1_.cpu_ready() && operand2_.cpu_ready());
285  }
286 
287  inline bool gpu_ready(int const deviceIndex) const { return false; }
288 
289  inline GPUWalkType gpu_init(IntVec const & extents,
290  GhostData const & ghosts,
291  IntVec const & hasBC,
292  int const deviceIndex,
293  cudaStream_t const & lhsStream) const {
294  return GPUWalkType(operand1_.gpu_init(extents,
295  ghosts,
296  hasBC,
297  deviceIndex,
298  lhsStream),
299  operand2_.gpu_init(extents,
300  ghosts,
301  hasBC,
302  deviceIndex,
303  lhsStream));
304  }
305 
306  inline void stream_wait_event(cudaEvent_t const & event) const {
307  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
308  }
309 
310  #ifdef NEBO_GPU_TEST
311  inline void gpu_prep(int const deviceIndex) const {
312  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
313  }
314  #endif
315  /* NEBO_GPU_TEST */
316  #endif
317  /* __CUDACC__ */
318 
319  private:
320  FunctorFactory const & functorFactory_;
321 
322  Operand1 const operand1_;
323 
324  Operand2 const operand2_;
325  };
326  #ifdef ENABLE_THREADS
327  template<typename FunctorFactory, typename Operand1, typename Operand2>
328  struct ApplyPointwise2<Resize, FunctorFactory, Operand1, Operand2> {
329  public:
330  ApplyPointwise2<SeqWalk,
331  FunctorFactory,
332  typename Operand1::SeqWalkType,
333  typename Operand2::SeqWalkType> typedef SeqWalkType
334  ;
335 
336  ApplyPointwise2(FunctorFactory const & functorFactory,
337  Operand1 const & operand1,
338  Operand2 const & operand2)
339  : functorFactory_(functorFactory), operand1_(operand1), operand2_(operand2)
340  {}
341 
342  inline SeqWalkType init(IntVec const & extents,
343  GhostData const & ghosts,
344  IntVec const & hasBC) const {
345  return SeqWalkType(functorFactory_,
346  operand1_.init(extents, ghosts, hasBC),
347  operand2_.init(extents, ghosts, hasBC));
348  }
349 
350  private:
351  FunctorFactory const & functorFactory_;
352 
353  Operand1 const operand1_;
354 
355  Operand2 const operand2_;
356  }
357  #endif
358  /* ENABLE_THREADS */;
359  template<typename FunctorFactory, typename Operand1, typename Operand2>
360  struct ApplyPointwise2<SeqWalk, FunctorFactory, Operand1, Operand2> {
361  public:
362  typename Operand1::value_type typedef value_type;
363 
364  ApplyPointwise2(FunctorFactory const & functorFactory,
365  Operand1 const & operand1,
366  Operand2 const & operand2)
367  : functor_(functorFactory()), operand1_(operand1), operand2_(operand2)
368  {}
369 
370  inline value_type eval(int const x, int const y, int const z) const {
371  return functor_(operand1_.eval(x, y, z), operand2_.eval(x, y, z));
372  }
373 
374  private:
375  typename FunctorFactory::Operator functor_;
376 
377  Operand1 operand1_;
378 
379  Operand2 operand2_;
380  };
381  #ifdef __CUDACC__
382  template<typename FunctorFactory, typename Operand1, typename Operand2>
383  struct ApplyPointwise2<GPUWalk, FunctorFactory, Operand1, Operand2> {
384  public:
385  typename Operand1::value_type typedef value_type;
386 
387  ApplyPointwise2(Operand1 const & operand1,
388  Operand2 const & operand2)
389  : operand1_(operand1), operand2_(operand2)
390  {}
391 
392  __device__ inline value_type eval(int const x,
393  int const y,
394  int const z) const { return 0; }
395 
396  private:
397  Operand1 operand1_;
398 
399  Operand2 operand2_;
400  }
401  #endif
402  /* __CUDACC__ */;
403 
404  template<typename Operand1, typename Operand2>
405  struct FindFieldType2 {
408  Result typedef Result;
409  };
410 
411  template<typename FunctorFactory, typename Operand1, typename Operand2>
412  inline NeboExpression<ApplyPointwise2<Initial,
413  FunctorFactory,
414  typename FinalType<Operand1>::
415  Result,
416  typename FinalType<Operand2>::
417  Result>,
419  apply_pointwise(FunctorFactory const & functorFactory,
420  Operand1 operand1,
421  Operand2 operand2) {
422  ApplyPointwise2<Initial,
423  FunctorFactory,
425  typename FinalType<Operand2>::Result> typedef
426  ReturnType;
427 
428  NeboExpression<ReturnType,
430  typedef ReturnTerm;
431 
432  return ReturnTerm(ReturnType(functorFactory,
433  normalize(operand1),
434  normalize(operand2)));
435  }
436 
437  template<typename CurrentMode,
438  typename FunctorFactory,
439  typename Operand1,
440  typename Operand2,
441  typename Operand3>
443  template<typename FunctorFactory,
444  typename Operand1,
445  typename Operand2,
446  typename Operand3>
447  struct ApplyPointwise3<Initial,
448  FunctorFactory,
449  Operand1,
450  Operand2,
451  Operand3> {
452  public:
453  ApplyPointwise3<SeqWalk,
454  FunctorFactory,
455  typename Operand1::SeqWalkType,
456  typename Operand2::SeqWalkType,
457  typename Operand3::SeqWalkType> typedef SeqWalkType;
458 
459  #ifdef ENABLE_THREADS
460  ApplyPointwise3<Resize,
461  FunctorFactory,
462  typename Operand1::ResizeType,
463  typename Operand2::ResizeType,
464  typename Operand3::ResizeType> typedef ResizeType;
465  #endif
466  /* ENABLE_THREADS */
467 
468  #ifdef __CUDACC__
469  ApplyPointwise3<GPUWalk,
470  FunctorFactory,
471  typename Operand1::GPUWalkType,
472  typename Operand2::GPUWalkType,
473  typename Operand3::GPUWalkType> typedef GPUWalkType
474  ;
475  #endif
476  /* __CUDACC__ */
477 
478  ApplyPointwise3(FunctorFactory const & functorFactory,
479  Operand1 const & operand1,
480  Operand2 const & operand2,
481  Operand3 const & operand3)
482  : functorFactory_(functorFactory),
483  operand1_(operand1),
484  operand2_(operand2),
485  operand3_(operand3)
486  {}
487 
488  inline GhostData ghosts_with_bc(void) const {
489  return min(operand1_.ghosts_with_bc(),
490  min(operand2_.ghosts_with_bc(), operand3_.ghosts_with_bc()));
491  }
492 
493  inline GhostData ghosts_without_bc(void) const {
494  return min(operand1_.ghosts_without_bc(),
495  min(operand2_.ghosts_without_bc(), operand3_.ghosts_without_bc()));
496  }
497 
498  inline bool has_extents(void) const {
499  return (operand1_.has_extents() || operand2_.has_extents() ||
500  operand3_.has_extents());
501  }
502 
503  inline IntVec extents(void) const {
504  #ifndef NDEBUG
505  if((operand1_.has_extents() || operand2_.has_extents() ||
506  operand3_.has_extents())) {
507  IntVec extents;
508 
509  if(operand1_.has_extents()) { extents = operand1_.extents(); }
510  else {
511  if(operand2_.has_extents()) {
512  extents = operand2_.extents();
513  }
514  else { extents = operand3_.extents(); };
515  };
516 
517  if(operand1_.has_extents()) {
518  assert(extents == operand1_.extents());
519  };
520 
521  if(operand2_.has_extents()) {
522  assert(extents == operand2_.extents());
523  };
524 
525  if(operand3_.has_extents()) {
526  assert(extents == operand3_.extents());
527  };
528  }
529  #endif
530  /* NDEBUG */;
531 
532  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
533  ?
534  operand2_.extents()
535  : (operand3_.extents()))));
536  }
537 
538  inline IntVec has_bc(void) const {
539  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
540  ?
541  operand2_.has_bc()
542  : (operand3_.has_bc()))));
543  }
544 
545  inline SeqWalkType init(IntVec const & extents,
546  GhostData const & ghosts,
547  IntVec const & hasBC) const {
548  return SeqWalkType(functorFactory_,
549  operand1_.init(extents, ghosts, hasBC),
550  operand2_.init(extents, ghosts, hasBC),
551  operand3_.init(extents, ghosts, hasBC));
552  }
553 
554  #ifdef ENABLE_THREADS
555  inline ResizeType resize(void) const {
556  return ResizeType(functorFactory_,
557  operand1_.resize(),
558  operand2_.resize(),
559  operand3_.resize());
560  }
561  #endif
562  /* ENABLE_THREADS */
563 
564  #ifdef __CUDACC__
565  inline bool cpu_ready(void) const {
566  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
567  operand3_.cpu_ready());
568  }
569 
570  inline bool gpu_ready(int const deviceIndex) const { return false; }
571 
572  inline GPUWalkType gpu_init(IntVec const & extents,
573  GhostData const & ghosts,
574  IntVec const & hasBC,
575  int const deviceIndex,
576  cudaStream_t const & lhsStream) const {
577  return GPUWalkType(operand1_.gpu_init(extents,
578  ghosts,
579  hasBC,
580  deviceIndex,
581  lhsStream),
582  operand2_.gpu_init(extents,
583  ghosts,
584  hasBC,
585  deviceIndex,
586  lhsStream),
587  operand3_.gpu_init(extents,
588  ghosts,
589  hasBC,
590  deviceIndex,
591  lhsStream));
592  }
593 
594  inline void stream_wait_event(cudaEvent_t const & event) const {
595  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
596  operand3_.stream_wait_event(event);
597  }
598 
599  #ifdef NEBO_GPU_TEST
600  inline void gpu_prep(int const deviceIndex) const {
601  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
602  operand3_.gpu_prep(deviceIndex);
603  }
604  #endif
605  /* NEBO_GPU_TEST */
606  #endif
607  /* __CUDACC__ */
608 
609  private:
610  FunctorFactory const & functorFactory_;
611 
612  Operand1 const operand1_;
613 
614  Operand2 const operand2_;
615 
616  Operand3 const operand3_;
617  };
618  #ifdef ENABLE_THREADS
619  template<typename FunctorFactory,
620  typename Operand1,
621  typename Operand2,
622  typename Operand3>
623  struct ApplyPointwise3<Resize,
624  FunctorFactory,
625  Operand1,
626  Operand2,
627  Operand3> {
628  public:
629  ApplyPointwise3<SeqWalk,
630  FunctorFactory,
631  typename Operand1::SeqWalkType,
632  typename Operand2::SeqWalkType,
633  typename Operand3::SeqWalkType> typedef SeqWalkType
634  ;
635 
636  ApplyPointwise3(FunctorFactory const & functorFactory,
637  Operand1 const & operand1,
638  Operand2 const & operand2,
639  Operand3 const & operand3)
640  : functorFactory_(functorFactory),
641  operand1_(operand1),
642  operand2_(operand2),
643  operand3_(operand3)
644  {}
645 
646  inline SeqWalkType init(IntVec const & extents,
647  GhostData const & ghosts,
648  IntVec const & hasBC) const {
649  return SeqWalkType(functorFactory_,
650  operand1_.init(extents, ghosts, hasBC),
651  operand2_.init(extents, ghosts, hasBC),
652  operand3_.init(extents, ghosts, hasBC));
653  }
654 
655  private:
656  FunctorFactory const & functorFactory_;
657 
658  Operand1 const operand1_;
659 
660  Operand2 const operand2_;
661 
662  Operand3 const operand3_;
663  }
664  #endif
665  /* ENABLE_THREADS */;
666  template<typename FunctorFactory,
667  typename Operand1,
668  typename Operand2,
669  typename Operand3>
670  struct ApplyPointwise3<SeqWalk,
671  FunctorFactory,
672  Operand1,
673  Operand2,
674  Operand3> {
675  public:
676  typename Operand1::value_type typedef value_type;
677 
678  ApplyPointwise3(FunctorFactory const & functorFactory,
679  Operand1 const & operand1,
680  Operand2 const & operand2,
681  Operand3 const & operand3)
682  : functor_(functorFactory()),
683  operand1_(operand1),
684  operand2_(operand2),
685  operand3_(operand3)
686  {}
687 
688  inline value_type eval(int const x, int const y, int const z) const {
689  return functor_(operand1_.eval(x, y, z),
690  operand2_.eval(x, y, z),
691  operand3_.eval(x, y, z));
692  }
693 
694  private:
695  typename FunctorFactory::Operator functor_;
696 
697  Operand1 operand1_;
698 
699  Operand2 operand2_;
700 
701  Operand3 operand3_;
702  };
703  #ifdef __CUDACC__
704  template<typename FunctorFactory,
705  typename Operand1,
706  typename Operand2,
707  typename Operand3>
708  struct ApplyPointwise3<GPUWalk,
709  FunctorFactory,
710  Operand1,
711  Operand2,
712  Operand3> {
713  public:
714  typename Operand1::value_type typedef value_type;
715 
716  ApplyPointwise3(Operand1 const & operand1,
717  Operand2 const & operand2,
718  Operand3 const & operand3)
719  : operand1_(operand1), operand2_(operand2), operand3_(operand3)
720  {}
721 
722  __device__ inline value_type eval(int const x,
723  int const y,
724  int const z) const { return 0; }
725 
726  private:
727  Operand1 operand1_;
728 
729  Operand2 operand2_;
730 
731  Operand3 operand3_;
732  }
733  #endif
734  /* __CUDACC__ */;
735 
736  template<typename Operand1, typename Operand2, typename Operand3>
737  struct FindFieldType3 {
739  typename RefineFieldType<typename
741  Result,
742  typename
744  Result>::Result>::
745  Result typedef Result;
746  };
747 
748  template<typename FunctorFactory,
749  typename Operand1,
750  typename Operand2,
751  typename Operand3>
752  inline NeboExpression<ApplyPointwise3<Initial,
753  FunctorFactory,
754  typename FinalType<Operand1>::
755  Result,
756  typename FinalType<Operand2>::
757  Result,
758  typename FinalType<Operand3>::
759  Result>,
760  typename FindFieldType3<Operand1,
761  Operand2,
762  Operand3>::Result>
763  apply_pointwise(FunctorFactory const & functorFactory,
764  Operand1 operand1,
765  Operand2 operand2,
766  Operand3 operand3) {
767  ApplyPointwise3<Initial,
768  FunctorFactory,
771  typename FinalType<Operand3>::Result> typedef
772  ReturnType;
773 
774  NeboExpression<ReturnType,
776  Result> typedef ReturnTerm;
777 
778  return ReturnTerm(ReturnType(functorFactory,
779  normalize(operand1),
780  normalize(operand2),
781  normalize(operand3)));
782  }
783 
784  template<typename CurrentMode,
785  typename FunctorFactory,
786  typename Operand1,
787  typename Operand2,
788  typename Operand3,
789  typename Operand4>
791  template<typename FunctorFactory,
792  typename Operand1,
793  typename Operand2,
794  typename Operand3,
795  typename Operand4>
796  struct ApplyPointwise4<Initial,
797  FunctorFactory,
798  Operand1,
799  Operand2,
800  Operand3,
801  Operand4> {
802  public:
803  ApplyPointwise4<SeqWalk,
804  FunctorFactory,
805  typename Operand1::SeqWalkType,
806  typename Operand2::SeqWalkType,
807  typename Operand3::SeqWalkType,
808  typename Operand4::SeqWalkType> typedef SeqWalkType;
809 
810  #ifdef ENABLE_THREADS
811  ApplyPointwise4<Resize,
812  FunctorFactory,
813  typename Operand1::ResizeType,
814  typename Operand2::ResizeType,
815  typename Operand3::ResizeType,
816  typename Operand4::ResizeType> typedef ResizeType;
817  #endif
818  /* ENABLE_THREADS */
819 
820  #ifdef __CUDACC__
821  ApplyPointwise4<GPUWalk,
822  FunctorFactory,
823  typename Operand1::GPUWalkType,
824  typename Operand2::GPUWalkType,
825  typename Operand3::GPUWalkType,
826  typename Operand4::GPUWalkType> typedef GPUWalkType
827  ;
828  #endif
829  /* __CUDACC__ */
830 
831  ApplyPointwise4(FunctorFactory const & functorFactory,
832  Operand1 const & operand1,
833  Operand2 const & operand2,
834  Operand3 const & operand3,
835  Operand4 const & operand4)
836  : functorFactory_(functorFactory),
837  operand1_(operand1),
838  operand2_(operand2),
839  operand3_(operand3),
840  operand4_(operand4)
841  {}
842 
843  inline GhostData ghosts_with_bc(void) const {
844  return min(operand1_.ghosts_with_bc(),
845  min(operand2_.ghosts_with_bc(),
846  min(operand3_.ghosts_with_bc(), operand4_.ghosts_with_bc())));
847  }
848 
849  inline GhostData ghosts_without_bc(void) const {
850  return min(operand1_.ghosts_without_bc(),
851  min(operand2_.ghosts_without_bc(),
852  min(operand3_.ghosts_without_bc(), operand4_.ghosts_without_bc())));
853  }
854 
855  inline bool has_extents(void) const {
856  return (operand1_.has_extents() || operand2_.has_extents() ||
857  operand3_.has_extents() || operand4_.has_extents());
858  }
859 
860  inline IntVec extents(void) const {
861  #ifndef NDEBUG
862  if((operand1_.has_extents() || operand2_.has_extents() ||
863  operand3_.has_extents() || operand4_.has_extents())) {
864  IntVec extents;
865 
866  if(operand1_.has_extents()) { extents = operand1_.extents(); }
867  else {
868  if(operand2_.has_extents()) {
869  extents = operand2_.extents();
870  }
871  else {
872  if(operand3_.has_extents()) {
873  extents = operand3_.extents();
874  }
875  else { extents = operand4_.extents(); };
876  };
877  };
878 
879  if(operand1_.has_extents()) {
880  assert(extents == operand1_.extents());
881  };
882 
883  if(operand2_.has_extents()) {
884  assert(extents == operand2_.extents());
885  };
886 
887  if(operand3_.has_extents()) {
888  assert(extents == operand3_.extents());
889  };
890 
891  if(operand4_.has_extents()) {
892  assert(extents == operand4_.extents());
893  };
894  }
895  #endif
896  /* NDEBUG */;
897 
898  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
899  ?
900  operand2_.extents()
901  : ((operand3_.has_extents()
902  ?
903  operand3_.extents()
904  : (operand4_.extents()))))));
905  }
906 
907  inline IntVec has_bc(void) const {
908  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
909  ?
910  operand2_.has_bc()
911  : ((operand3_.has_extents()
912  ?
913  operand3_.has_bc()
914  : (operand4_.has_bc()))))));
915  }
916 
917  inline SeqWalkType init(IntVec const & extents,
918  GhostData const & ghosts,
919  IntVec const & hasBC) const {
920  return SeqWalkType(functorFactory_,
921  operand1_.init(extents, ghosts, hasBC),
922  operand2_.init(extents, ghosts, hasBC),
923  operand3_.init(extents, ghosts, hasBC),
924  operand4_.init(extents, ghosts, hasBC));
925  }
926 
927  #ifdef ENABLE_THREADS
928  inline ResizeType resize(void) const {
929  return ResizeType(functorFactory_,
930  operand1_.resize(),
931  operand2_.resize(),
932  operand3_.resize(),
933  operand4_.resize());
934  }
935  #endif
936  /* ENABLE_THREADS */
937 
938  #ifdef __CUDACC__
939  inline bool cpu_ready(void) const {
940  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
941  operand3_.cpu_ready() && operand4_.cpu_ready());
942  }
943 
944  inline bool gpu_ready(int const deviceIndex) const { return false; }
945 
946  inline GPUWalkType gpu_init(IntVec const & extents,
947  GhostData const & ghosts,
948  IntVec const & hasBC,
949  int const deviceIndex,
950  cudaStream_t const & lhsStream) const {
951  return GPUWalkType(operand1_.gpu_init(extents,
952  ghosts,
953  hasBC,
954  deviceIndex,
955  lhsStream),
956  operand2_.gpu_init(extents,
957  ghosts,
958  hasBC,
959  deviceIndex,
960  lhsStream),
961  operand3_.gpu_init(extents,
962  ghosts,
963  hasBC,
964  deviceIndex,
965  lhsStream),
966  operand4_.gpu_init(extents,
967  ghosts,
968  hasBC,
969  deviceIndex,
970  lhsStream));
971  }
972 
973  inline void stream_wait_event(cudaEvent_t const & event) const {
974  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
975  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
976  }
977 
978  #ifdef NEBO_GPU_TEST
979  inline void gpu_prep(int const deviceIndex) const {
980  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
981  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
982  }
983  #endif
984  /* NEBO_GPU_TEST */
985  #endif
986  /* __CUDACC__ */
987 
988  private:
989  FunctorFactory const & functorFactory_;
990 
991  Operand1 const operand1_;
992 
993  Operand2 const operand2_;
994 
995  Operand3 const operand3_;
996 
997  Operand4 const operand4_;
998  };
999  #ifdef ENABLE_THREADS
1000  template<typename FunctorFactory,
1001  typename Operand1,
1002  typename Operand2,
1003  typename Operand3,
1004  typename Operand4>
1005  struct ApplyPointwise4<Resize,
1006  FunctorFactory,
1007  Operand1,
1008  Operand2,
1009  Operand3,
1010  Operand4> {
1011  public:
1012  ApplyPointwise4<SeqWalk,
1013  FunctorFactory,
1014  typename Operand1::SeqWalkType,
1015  typename Operand2::SeqWalkType,
1016  typename Operand3::SeqWalkType,
1017  typename Operand4::SeqWalkType> typedef SeqWalkType
1018  ;
1019 
1020  ApplyPointwise4(FunctorFactory const & functorFactory,
1021  Operand1 const & operand1,
1022  Operand2 const & operand2,
1023  Operand3 const & operand3,
1024  Operand4 const & operand4)
1025  : functorFactory_(functorFactory),
1026  operand1_(operand1),
1027  operand2_(operand2),
1028  operand3_(operand3),
1029  operand4_(operand4)
1030  {}
1031 
1032  inline SeqWalkType init(IntVec const & extents,
1033  GhostData const & ghosts,
1034  IntVec const & hasBC) const {
1035  return SeqWalkType(functorFactory_,
1036  operand1_.init(extents, ghosts, hasBC),
1037  operand2_.init(extents, ghosts, hasBC),
1038  operand3_.init(extents, ghosts, hasBC),
1039  operand4_.init(extents, ghosts, hasBC));
1040  }
1041 
1042  private:
1043  FunctorFactory const & functorFactory_;
1044 
1045  Operand1 const operand1_;
1046 
1047  Operand2 const operand2_;
1048 
1049  Operand3 const operand3_;
1050 
1051  Operand4 const operand4_;
1052  }
1053  #endif
1054  /* ENABLE_THREADS */;
1055  template<typename FunctorFactory,
1056  typename Operand1,
1057  typename Operand2,
1058  typename Operand3,
1059  typename Operand4>
1060  struct ApplyPointwise4<SeqWalk,
1061  FunctorFactory,
1062  Operand1,
1063  Operand2,
1064  Operand3,
1065  Operand4> {
1066  public:
1067  typename Operand1::value_type typedef value_type;
1068 
1069  ApplyPointwise4(FunctorFactory const & functorFactory,
1070  Operand1 const & operand1,
1071  Operand2 const & operand2,
1072  Operand3 const & operand3,
1073  Operand4 const & operand4)
1074  : functor_(functorFactory()),
1075  operand1_(operand1),
1076  operand2_(operand2),
1077  operand3_(operand3),
1078  operand4_(operand4)
1079  {}
1080 
1081  inline value_type eval(int const x, int const y, int const z) const {
1082  return functor_(operand1_.eval(x, y, z),
1083  operand2_.eval(x, y, z),
1084  operand3_.eval(x, y, z),
1085  operand4_.eval(x, y, z));
1086  }
1087 
1088  private:
1089  typename FunctorFactory::Operator functor_;
1090 
1091  Operand1 operand1_;
1092 
1093  Operand2 operand2_;
1094 
1095  Operand3 operand3_;
1096 
1097  Operand4 operand4_;
1098  };
1099  #ifdef __CUDACC__
1100  template<typename FunctorFactory,
1101  typename Operand1,
1102  typename Operand2,
1103  typename Operand3,
1104  typename Operand4>
1105  struct ApplyPointwise4<GPUWalk,
1106  FunctorFactory,
1107  Operand1,
1108  Operand2,
1109  Operand3,
1110  Operand4> {
1111  public:
1112  typename Operand1::value_type typedef value_type;
1113 
1114  ApplyPointwise4(Operand1 const & operand1,
1115  Operand2 const & operand2,
1116  Operand3 const & operand3,
1117  Operand4 const & operand4)
1118  : operand1_(operand1),
1119  operand2_(operand2),
1120  operand3_(operand3),
1121  operand4_(operand4)
1122  {}
1123 
1124  __device__ inline value_type eval(int const x,
1125  int const y,
1126  int const z) const { return 0; }
1127 
1128  private:
1129  Operand1 operand1_;
1130 
1131  Operand2 operand2_;
1132 
1133  Operand3 operand3_;
1134 
1135  Operand4 operand4_;
1136  }
1137  #endif
1138  /* __CUDACC__ */;
1139 
1140  template<typename Operand1,
1141  typename Operand2,
1142  typename Operand3,
1143  typename Operand4>
1146  typename RefineFieldType<typename
1148  Result,
1149  typename
1150  RefineFieldType<typename
1152  Result,
1153  typename
1155  Result>::
1156  Result>::Result>::
1157  Result typedef Result;
1158  };
1159 
1160  template<typename FunctorFactory,
1161  typename Operand1,
1162  typename Operand2,
1163  typename Operand3,
1164  typename Operand4>
1165  inline NeboExpression<ApplyPointwise4<Initial,
1166  FunctorFactory,
1167  typename FinalType<Operand1>::
1168  Result,
1169  typename FinalType<Operand2>::
1170  Result,
1171  typename FinalType<Operand3>::
1172  Result,
1173  typename FinalType<Operand4>::
1174  Result>,
1175  typename FindFieldType4<Operand1,
1176  Operand2,
1177  Operand3,
1178  Operand4>::Result>
1179  apply_pointwise(FunctorFactory const & functorFactory,
1180  Operand1 operand1,
1181  Operand2 operand2,
1182  Operand3 operand3,
1183  Operand4 operand4) {
1184  ApplyPointwise4<Initial,
1185  FunctorFactory,
1186  typename FinalType<Operand1>::Result,
1187  typename FinalType<Operand2>::Result,
1188  typename FinalType<Operand3>::Result,
1189  typename FinalType<Operand4>::Result> typedef
1190  ReturnType;
1191 
1192  NeboExpression<ReturnType,
1193  typename FindFieldType4<Operand1,
1194  Operand2,
1195  Operand3,
1196  Operand4>::Result> typedef
1197  ReturnTerm;
1198 
1199  return ReturnTerm(ReturnType(functorFactory,
1200  normalize(operand1),
1201  normalize(operand2),
1202  normalize(operand3),
1203  normalize(operand4)));
1204  }
1205 
1206  template<typename CurrentMode,
1207  typename FunctorFactory,
1208  typename Operand1,
1209  typename Operand2,
1210  typename Operand3,
1211  typename Operand4,
1212  typename Operand5>
1214  template<typename FunctorFactory,
1215  typename Operand1,
1216  typename Operand2,
1217  typename Operand3,
1218  typename Operand4,
1219  typename Operand5>
1220  struct ApplyPointwise5<Initial,
1221  FunctorFactory,
1222  Operand1,
1223  Operand2,
1224  Operand3,
1225  Operand4,
1226  Operand5> {
1227  public:
1228  ApplyPointwise5<SeqWalk,
1229  FunctorFactory,
1230  typename Operand1::SeqWalkType,
1231  typename Operand2::SeqWalkType,
1232  typename Operand3::SeqWalkType,
1233  typename Operand4::SeqWalkType,
1234  typename Operand5::SeqWalkType> typedef SeqWalkType;
1235 
1236  #ifdef ENABLE_THREADS
1237  ApplyPointwise5<Resize,
1238  FunctorFactory,
1239  typename Operand1::ResizeType,
1240  typename Operand2::ResizeType,
1241  typename Operand3::ResizeType,
1242  typename Operand4::ResizeType,
1243  typename Operand5::ResizeType> typedef ResizeType;
1244  #endif
1245  /* ENABLE_THREADS */
1246 
1247  #ifdef __CUDACC__
1248  ApplyPointwise5<GPUWalk,
1249  FunctorFactory,
1250  typename Operand1::GPUWalkType,
1251  typename Operand2::GPUWalkType,
1252  typename Operand3::GPUWalkType,
1253  typename Operand4::GPUWalkType,
1254  typename Operand5::GPUWalkType> typedef GPUWalkType
1255  ;
1256  #endif
1257  /* __CUDACC__ */
1258 
1259  ApplyPointwise5(FunctorFactory const & functorFactory,
1260  Operand1 const & operand1,
1261  Operand2 const & operand2,
1262  Operand3 const & operand3,
1263  Operand4 const & operand4,
1264  Operand5 const & operand5)
1265  : functorFactory_(functorFactory),
1266  operand1_(operand1),
1267  operand2_(operand2),
1268  operand3_(operand3),
1269  operand4_(operand4),
1270  operand5_(operand5)
1271  {}
1272 
1273  inline GhostData ghosts_with_bc(void) const {
1274  return min(operand1_.ghosts_with_bc(),
1275  min(operand2_.ghosts_with_bc(),
1276  min(operand3_.ghosts_with_bc(),
1277  min(operand4_.ghosts_with_bc(), operand5_.ghosts_with_bc()))));
1278  }
1279 
1280  inline GhostData ghosts_without_bc(void) const {
1281  return min(operand1_.ghosts_without_bc(),
1282  min(operand2_.ghosts_without_bc(),
1283  min(operand3_.ghosts_without_bc(),
1284  min(operand4_.ghosts_without_bc(), operand5_.ghosts_without_bc()))));
1285  }
1286 
1287  inline bool has_extents(void) const {
1288  return (operand1_.has_extents() || operand2_.has_extents() ||
1289  operand3_.has_extents() || operand4_.has_extents() ||
1290  operand5_.has_extents());
1291  }
1292 
1293  inline IntVec extents(void) const {
1294  #ifndef NDEBUG
1295  if((operand1_.has_extents() || operand2_.has_extents() ||
1296  operand3_.has_extents() || operand4_.has_extents() ||
1297  operand5_.has_extents())) {
1298  IntVec extents;
1299 
1300  if(operand1_.has_extents()) { extents = operand1_.extents(); }
1301  else {
1302  if(operand2_.has_extents()) {
1303  extents = operand2_.extents();
1304  }
1305  else {
1306  if(operand3_.has_extents()) {
1307  extents = operand3_.extents();
1308  }
1309  else {
1310  if(operand4_.has_extents()) {
1311  extents = operand4_.extents();
1312  }
1313  else { extents = operand5_.extents(); };
1314  };
1315  };
1316  };
1317 
1318  if(operand1_.has_extents()) {
1319  assert(extents == operand1_.extents());
1320  };
1321 
1322  if(operand2_.has_extents()) {
1323  assert(extents == operand2_.extents());
1324  };
1325 
1326  if(operand3_.has_extents()) {
1327  assert(extents == operand3_.extents());
1328  };
1329 
1330  if(operand4_.has_extents()) {
1331  assert(extents == operand4_.extents());
1332  };
1333 
1334  if(operand5_.has_extents()) {
1335  assert(extents == operand5_.extents());
1336  };
1337  }
1338  #endif
1339  /* NDEBUG */;
1340 
1341  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
1342  ?
1343  operand2_.extents()
1344  : ((operand3_.has_extents()
1345  ?
1346  operand3_.extents()
1347  : ((operand4_.has_extents()
1348  ?
1349  operand4_.extents()
1350  :
1351  (operand5_.extents()))))))));
1352  }
1353 
1354  inline IntVec has_bc(void) const {
1355  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
1356  ?
1357  operand2_.has_bc()
1358  : ((operand3_.has_extents()
1359  ?
1360  operand3_.has_bc()
1361  : ((operand4_.has_extents()
1362  ?
1363  operand4_.has_bc()
1364  : (operand5_.has_bc()))))))));
1365  }
1366 
1367  inline SeqWalkType init(IntVec const & extents,
1368  GhostData const & ghosts,
1369  IntVec const & hasBC) const {
1370  return SeqWalkType(functorFactory_,
1371  operand1_.init(extents, ghosts, hasBC),
1372  operand2_.init(extents, ghosts, hasBC),
1373  operand3_.init(extents, ghosts, hasBC),
1374  operand4_.init(extents, ghosts, hasBC),
1375  operand5_.init(extents, ghosts, hasBC));
1376  }
1377 
1378  #ifdef ENABLE_THREADS
1379  inline ResizeType resize(void) const {
1380  return ResizeType(functorFactory_,
1381  operand1_.resize(),
1382  operand2_.resize(),
1383  operand3_.resize(),
1384  operand4_.resize(),
1385  operand5_.resize());
1386  }
1387  #endif
1388  /* ENABLE_THREADS */
1389 
1390  #ifdef __CUDACC__
1391  inline bool cpu_ready(void) const {
1392  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
1393  operand3_.cpu_ready() && operand4_.cpu_ready() &&
1394  operand5_.cpu_ready());
1395  }
1396 
1397  inline bool gpu_ready(int const deviceIndex) const { return false; }
1398 
1399  inline GPUWalkType gpu_init(IntVec const & extents,
1400  GhostData const & ghosts,
1401  IntVec const & hasBC,
1402  int const deviceIndex,
1403  cudaStream_t const & lhsStream) const {
1404  return GPUWalkType(operand1_.gpu_init(extents,
1405  ghosts,
1406  hasBC,
1407  deviceIndex,
1408  lhsStream),
1409  operand2_.gpu_init(extents,
1410  ghosts,
1411  hasBC,
1412  deviceIndex,
1413  lhsStream),
1414  operand3_.gpu_init(extents,
1415  ghosts,
1416  hasBC,
1417  deviceIndex,
1418  lhsStream),
1419  operand4_.gpu_init(extents,
1420  ghosts,
1421  hasBC,
1422  deviceIndex,
1423  lhsStream),
1424  operand5_.gpu_init(extents,
1425  ghosts,
1426  hasBC,
1427  deviceIndex,
1428  lhsStream));
1429  }
1430 
1431  inline void stream_wait_event(cudaEvent_t const & event) const {
1432  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
1433  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
1434  operand5_.stream_wait_event(event);
1435  }
1436 
1437  #ifdef NEBO_GPU_TEST
1438  inline void gpu_prep(int const deviceIndex) const {
1439  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
1440  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
1441  operand5_.gpu_prep(deviceIndex);
1442  }
1443  #endif
1444  /* NEBO_GPU_TEST */
1445  #endif
1446  /* __CUDACC__ */
1447 
1448  private:
1449  FunctorFactory const & functorFactory_;
1450 
1451  Operand1 const operand1_;
1452 
1453  Operand2 const operand2_;
1454 
1455  Operand3 const operand3_;
1456 
1457  Operand4 const operand4_;
1458 
1459  Operand5 const operand5_;
1460  };
1461  #ifdef ENABLE_THREADS
1462  template<typename FunctorFactory,
1463  typename Operand1,
1464  typename Operand2,
1465  typename Operand3,
1466  typename Operand4,
1467  typename Operand5>
1468  struct ApplyPointwise5<Resize,
1469  FunctorFactory,
1470  Operand1,
1471  Operand2,
1472  Operand3,
1473  Operand4,
1474  Operand5> {
1475  public:
1476  ApplyPointwise5<SeqWalk,
1477  FunctorFactory,
1478  typename Operand1::SeqWalkType,
1479  typename Operand2::SeqWalkType,
1480  typename Operand3::SeqWalkType,
1481  typename Operand4::SeqWalkType,
1482  typename Operand5::SeqWalkType> typedef SeqWalkType
1483  ;
1484 
1485  ApplyPointwise5(FunctorFactory const & functorFactory,
1486  Operand1 const & operand1,
1487  Operand2 const & operand2,
1488  Operand3 const & operand3,
1489  Operand4 const & operand4,
1490  Operand5 const & operand5)
1491  : functorFactory_(functorFactory),
1492  operand1_(operand1),
1493  operand2_(operand2),
1494  operand3_(operand3),
1495  operand4_(operand4),
1496  operand5_(operand5)
1497  {}
1498 
1499  inline SeqWalkType init(IntVec const & extents,
1500  GhostData const & ghosts,
1501  IntVec const & hasBC) const {
1502  return SeqWalkType(functorFactory_,
1503  operand1_.init(extents, ghosts, hasBC),
1504  operand2_.init(extents, ghosts, hasBC),
1505  operand3_.init(extents, ghosts, hasBC),
1506  operand4_.init(extents, ghosts, hasBC),
1507  operand5_.init(extents, ghosts, hasBC));
1508  }
1509 
1510  private:
1511  FunctorFactory const & functorFactory_;
1512 
1513  Operand1 const operand1_;
1514 
1515  Operand2 const operand2_;
1516 
1517  Operand3 const operand3_;
1518 
1519  Operand4 const operand4_;
1520 
1521  Operand5 const operand5_;
1522  }
1523  #endif
1524  /* ENABLE_THREADS */;
1525  template<typename FunctorFactory,
1526  typename Operand1,
1527  typename Operand2,
1528  typename Operand3,
1529  typename Operand4,
1530  typename Operand5>
1531  struct ApplyPointwise5<SeqWalk,
1532  FunctorFactory,
1533  Operand1,
1534  Operand2,
1535  Operand3,
1536  Operand4,
1537  Operand5> {
1538  public:
1539  typename Operand1::value_type typedef value_type;
1540 
1541  ApplyPointwise5(FunctorFactory const & functorFactory,
1542  Operand1 const & operand1,
1543  Operand2 const & operand2,
1544  Operand3 const & operand3,
1545  Operand4 const & operand4,
1546  Operand5 const & operand5)
1547  : functor_(functorFactory()),
1548  operand1_(operand1),
1549  operand2_(operand2),
1550  operand3_(operand3),
1551  operand4_(operand4),
1552  operand5_(operand5)
1553  {}
1554 
1555  inline value_type eval(int const x, int const y, int const z) const {
1556  return functor_(operand1_.eval(x, y, z),
1557  operand2_.eval(x, y, z),
1558  operand3_.eval(x, y, z),
1559  operand4_.eval(x, y, z),
1560  operand5_.eval(x, y, z));
1561  }
1562 
1563  private:
1564  typename FunctorFactory::Operator functor_;
1565 
1566  Operand1 operand1_;
1567 
1568  Operand2 operand2_;
1569 
1570  Operand3 operand3_;
1571 
1572  Operand4 operand4_;
1573 
1574  Operand5 operand5_;
1575  };
1576  #ifdef __CUDACC__
1577  template<typename FunctorFactory,
1578  typename Operand1,
1579  typename Operand2,
1580  typename Operand3,
1581  typename Operand4,
1582  typename Operand5>
1583  struct ApplyPointwise5<GPUWalk,
1584  FunctorFactory,
1585  Operand1,
1586  Operand2,
1587  Operand3,
1588  Operand4,
1589  Operand5> {
1590  public:
1591  typename Operand1::value_type typedef value_type;
1592 
1593  ApplyPointwise5(Operand1 const & operand1,
1594  Operand2 const & operand2,
1595  Operand3 const & operand3,
1596  Operand4 const & operand4,
1597  Operand5 const & operand5)
1598  : operand1_(operand1),
1599  operand2_(operand2),
1600  operand3_(operand3),
1601  operand4_(operand4),
1602  operand5_(operand5)
1603  {}
1604 
1605  __device__ inline value_type eval(int const x,
1606  int const y,
1607  int const z) const { return 0; }
1608 
1609  private:
1610  Operand1 operand1_;
1611 
1612  Operand2 operand2_;
1613 
1614  Operand3 operand3_;
1615 
1616  Operand4 operand4_;
1617 
1618  Operand5 operand5_;
1619  }
1620  #endif
1621  /* __CUDACC__ */;
1622 
1623  template<typename Operand1,
1624  typename Operand2,
1625  typename Operand3,
1626  typename Operand4,
1627  typename Operand5>
1630  typename RefineFieldType<typename
1632  Result,
1633  typename
1634  RefineFieldType<typename
1636  Result,
1637  typename
1638  RefineFieldType<typename
1640  Result,
1641  typename
1643  Result>::
1644  Result>::
1645  Result>::Result>::
1646  Result typedef Result;
1647  };
1648 
1649  template<typename FunctorFactory,
1650  typename Operand1,
1651  typename Operand2,
1652  typename Operand3,
1653  typename Operand4,
1654  typename Operand5>
1655  inline NeboExpression<ApplyPointwise5<Initial,
1656  FunctorFactory,
1657  typename FinalType<Operand1>::
1658  Result,
1659  typename FinalType<Operand2>::
1660  Result,
1661  typename FinalType<Operand3>::
1662  Result,
1663  typename FinalType<Operand4>::
1664  Result,
1665  typename FinalType<Operand5>::
1666  Result>,
1667  typename FindFieldType5<Operand1,
1668  Operand2,
1669  Operand3,
1670  Operand4,
1671  Operand5>::Result>
1672  apply_pointwise(FunctorFactory const & functorFactory,
1673  Operand1 operand1,
1674  Operand2 operand2,
1675  Operand3 operand3,
1676  Operand4 operand4,
1677  Operand5 operand5) {
1678  ApplyPointwise5<Initial,
1679  FunctorFactory,
1680  typename FinalType<Operand1>::Result,
1681  typename FinalType<Operand2>::Result,
1682  typename FinalType<Operand3>::Result,
1683  typename FinalType<Operand4>::Result,
1684  typename FinalType<Operand5>::Result> typedef
1685  ReturnType;
1686 
1687  NeboExpression<ReturnType,
1688  typename FindFieldType5<Operand1,
1689  Operand2,
1690  Operand3,
1691  Operand4,
1692  Operand5>::Result> typedef
1693  ReturnTerm;
1694 
1695  return ReturnTerm(ReturnType(functorFactory,
1696  normalize(operand1),
1697  normalize(operand2),
1698  normalize(operand3),
1699  normalize(operand4),
1700  normalize(operand5)));
1701  }
1702 
1703  template<typename CurrentMode,
1704  typename FunctorFactory,
1705  typename Operand1,
1706  typename Operand2,
1707  typename Operand3,
1708  typename Operand4,
1709  typename Operand5,
1710  typename Operand6>
1712  template<typename FunctorFactory,
1713  typename Operand1,
1714  typename Operand2,
1715  typename Operand3,
1716  typename Operand4,
1717  typename Operand5,
1718  typename Operand6>
1719  struct ApplyPointwise6<Initial,
1720  FunctorFactory,
1721  Operand1,
1722  Operand2,
1723  Operand3,
1724  Operand4,
1725  Operand5,
1726  Operand6> {
1727  public:
1728  ApplyPointwise6<SeqWalk,
1729  FunctorFactory,
1730  typename Operand1::SeqWalkType,
1731  typename Operand2::SeqWalkType,
1732  typename Operand3::SeqWalkType,
1733  typename Operand4::SeqWalkType,
1734  typename Operand5::SeqWalkType,
1735  typename Operand6::SeqWalkType> typedef SeqWalkType;
1736 
1737  #ifdef ENABLE_THREADS
1738  ApplyPointwise6<Resize,
1739  FunctorFactory,
1740  typename Operand1::ResizeType,
1741  typename Operand2::ResizeType,
1742  typename Operand3::ResizeType,
1743  typename Operand4::ResizeType,
1744  typename Operand5::ResizeType,
1745  typename Operand6::ResizeType> typedef ResizeType;
1746  #endif
1747  /* ENABLE_THREADS */
1748 
1749  #ifdef __CUDACC__
1750  ApplyPointwise6<GPUWalk,
1751  FunctorFactory,
1752  typename Operand1::GPUWalkType,
1753  typename Operand2::GPUWalkType,
1754  typename Operand3::GPUWalkType,
1755  typename Operand4::GPUWalkType,
1756  typename Operand5::GPUWalkType,
1757  typename Operand6::GPUWalkType> typedef GPUWalkType
1758  ;
1759  #endif
1760  /* __CUDACC__ */
1761 
1762  ApplyPointwise6(FunctorFactory const & functorFactory,
1763  Operand1 const & operand1,
1764  Operand2 const & operand2,
1765  Operand3 const & operand3,
1766  Operand4 const & operand4,
1767  Operand5 const & operand5,
1768  Operand6 const & operand6)
1769  : functorFactory_(functorFactory),
1770  operand1_(operand1),
1771  operand2_(operand2),
1772  operand3_(operand3),
1773  operand4_(operand4),
1774  operand5_(operand5),
1775  operand6_(operand6)
1776  {}
1777 
1778  inline GhostData ghosts_with_bc(void) const {
1779  return min(operand1_.ghosts_with_bc(),
1780  min(operand2_.ghosts_with_bc(),
1781  min(operand3_.ghosts_with_bc(),
1782  min(operand4_.ghosts_with_bc(),
1783  min(operand5_.ghosts_with_bc(), operand6_.ghosts_with_bc())))));
1784  }
1785 
1786  inline GhostData ghosts_without_bc(void) const {
1787  return min(operand1_.ghosts_without_bc(),
1788  min(operand2_.ghosts_without_bc(),
1789  min(operand3_.ghosts_without_bc(),
1790  min(operand4_.ghosts_without_bc(),
1791  min(operand5_.ghosts_without_bc(), operand6_.ghosts_without_bc())))));
1792  }
1793 
1794  inline bool has_extents(void) const {
1795  return (operand1_.has_extents() || operand2_.has_extents() ||
1796  operand3_.has_extents() || operand4_.has_extents() ||
1797  operand5_.has_extents() || operand6_.has_extents());
1798  }
1799 
1800  inline IntVec extents(void) const {
1801  #ifndef NDEBUG
1802  if((operand1_.has_extents() || operand2_.has_extents() ||
1803  operand3_.has_extents() || operand4_.has_extents() ||
1804  operand5_.has_extents() || operand6_.has_extents())) {
1805  IntVec extents;
1806 
1807  if(operand1_.has_extents()) { extents = operand1_.extents(); }
1808  else {
1809  if(operand2_.has_extents()) {
1810  extents = operand2_.extents();
1811  }
1812  else {
1813  if(operand3_.has_extents()) {
1814  extents = operand3_.extents();
1815  }
1816  else {
1817  if(operand4_.has_extents()) {
1818  extents = operand4_.extents();
1819  }
1820  else {
1821  if(operand5_.has_extents()) {
1822  extents = operand5_.extents();
1823  }
1824  else { extents = operand6_.extents(); };
1825  };
1826  };
1827  };
1828  };
1829 
1830  if(operand1_.has_extents()) {
1831  assert(extents == operand1_.extents());
1832  };
1833 
1834  if(operand2_.has_extents()) {
1835  assert(extents == operand2_.extents());
1836  };
1837 
1838  if(operand3_.has_extents()) {
1839  assert(extents == operand3_.extents());
1840  };
1841 
1842  if(operand4_.has_extents()) {
1843  assert(extents == operand4_.extents());
1844  };
1845 
1846  if(operand5_.has_extents()) {
1847  assert(extents == operand5_.extents());
1848  };
1849 
1850  if(operand6_.has_extents()) {
1851  assert(extents == operand6_.extents());
1852  };
1853  }
1854  #endif
1855  /* NDEBUG */;
1856 
1857  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
1858  ?
1859  operand2_.extents()
1860  : ((operand3_.has_extents()
1861  ?
1862  operand3_.extents()
1863  : ((operand4_.has_extents()
1864  ?
1865  operand4_.extents()
1866  :
1867  ((operand5_.has_extents()
1868  ?
1869  operand5_.extents()
1870  :
1871  (operand6_.extents()))))))))));
1872  }
1873 
1874  inline IntVec has_bc(void) const {
1875  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
1876  ?
1877  operand2_.has_bc()
1878  : ((operand3_.has_extents()
1879  ?
1880  operand3_.has_bc()
1881  : ((operand4_.has_extents()
1882  ?
1883  operand4_.has_bc()
1884  : ((operand5_.has_extents()
1885  ?
1886  operand5_.has_bc()
1887  :
1888  (operand6_.has_bc()))))))))));
1889  }
1890 
1891  inline SeqWalkType init(IntVec const & extents,
1892  GhostData const & ghosts,
1893  IntVec const & hasBC) const {
1894  return SeqWalkType(functorFactory_,
1895  operand1_.init(extents, ghosts, hasBC),
1896  operand2_.init(extents, ghosts, hasBC),
1897  operand3_.init(extents, ghosts, hasBC),
1898  operand4_.init(extents, ghosts, hasBC),
1899  operand5_.init(extents, ghosts, hasBC),
1900  operand6_.init(extents, ghosts, hasBC));
1901  }
1902 
1903  #ifdef ENABLE_THREADS
1904  inline ResizeType resize(void) const {
1905  return ResizeType(functorFactory_,
1906  operand1_.resize(),
1907  operand2_.resize(),
1908  operand3_.resize(),
1909  operand4_.resize(),
1910  operand5_.resize(),
1911  operand6_.resize());
1912  }
1913  #endif
1914  /* ENABLE_THREADS */
1915 
1916  #ifdef __CUDACC__
1917  inline bool cpu_ready(void) const {
1918  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
1919  operand3_.cpu_ready() && operand4_.cpu_ready() &&
1920  operand5_.cpu_ready() && operand6_.cpu_ready());
1921  }
1922 
1923  inline bool gpu_ready(int const deviceIndex) const { return false; }
1924 
1925  inline GPUWalkType gpu_init(IntVec const & extents,
1926  GhostData const & ghosts,
1927  IntVec const & hasBC,
1928  int const deviceIndex,
1929  cudaStream_t const & lhsStream) const {
1930  return GPUWalkType(operand1_.gpu_init(extents,
1931  ghosts,
1932  hasBC,
1933  deviceIndex,
1934  lhsStream),
1935  operand2_.gpu_init(extents,
1936  ghosts,
1937  hasBC,
1938  deviceIndex,
1939  lhsStream),
1940  operand3_.gpu_init(extents,
1941  ghosts,
1942  hasBC,
1943  deviceIndex,
1944  lhsStream),
1945  operand4_.gpu_init(extents,
1946  ghosts,
1947  hasBC,
1948  deviceIndex,
1949  lhsStream),
1950  operand5_.gpu_init(extents,
1951  ghosts,
1952  hasBC,
1953  deviceIndex,
1954  lhsStream),
1955  operand6_.gpu_init(extents,
1956  ghosts,
1957  hasBC,
1958  deviceIndex,
1959  lhsStream));
1960  }
1961 
1962  inline void stream_wait_event(cudaEvent_t const & event) const {
1963  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
1964  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
1965  operand5_.stream_wait_event(event); operand6_.stream_wait_event(event);
1966  }
1967 
1968  #ifdef NEBO_GPU_TEST
1969  inline void gpu_prep(int const deviceIndex) const {
1970  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
1971  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
1972  operand5_.gpu_prep(deviceIndex); operand6_.gpu_prep(deviceIndex);
1973  }
1974  #endif
1975  /* NEBO_GPU_TEST */
1976  #endif
1977  /* __CUDACC__ */
1978 
1979  private:
1980  FunctorFactory const & functorFactory_;
1981 
1982  Operand1 const operand1_;
1983 
1984  Operand2 const operand2_;
1985 
1986  Operand3 const operand3_;
1987 
1988  Operand4 const operand4_;
1989 
1990  Operand5 const operand5_;
1991 
1992  Operand6 const operand6_;
1993  };
1994  #ifdef ENABLE_THREADS
1995  template<typename FunctorFactory,
1996  typename Operand1,
1997  typename Operand2,
1998  typename Operand3,
1999  typename Operand4,
2000  typename Operand5,
2001  typename Operand6>
2002  struct ApplyPointwise6<Resize,
2003  FunctorFactory,
2004  Operand1,
2005  Operand2,
2006  Operand3,
2007  Operand4,
2008  Operand5,
2009  Operand6> {
2010  public:
2011  ApplyPointwise6<SeqWalk,
2012  FunctorFactory,
2013  typename Operand1::SeqWalkType,
2014  typename Operand2::SeqWalkType,
2015  typename Operand3::SeqWalkType,
2016  typename Operand4::SeqWalkType,
2017  typename Operand5::SeqWalkType,
2018  typename Operand6::SeqWalkType> typedef SeqWalkType
2019  ;
2020 
2021  ApplyPointwise6(FunctorFactory const & functorFactory,
2022  Operand1 const & operand1,
2023  Operand2 const & operand2,
2024  Operand3 const & operand3,
2025  Operand4 const & operand4,
2026  Operand5 const & operand5,
2027  Operand6 const & operand6)
2028  : functorFactory_(functorFactory),
2029  operand1_(operand1),
2030  operand2_(operand2),
2031  operand3_(operand3),
2032  operand4_(operand4),
2033  operand5_(operand5),
2034  operand6_(operand6)
2035  {}
2036 
2037  inline SeqWalkType init(IntVec const & extents,
2038  GhostData const & ghosts,
2039  IntVec const & hasBC) const {
2040  return SeqWalkType(functorFactory_,
2041  operand1_.init(extents, ghosts, hasBC),
2042  operand2_.init(extents, ghosts, hasBC),
2043  operand3_.init(extents, ghosts, hasBC),
2044  operand4_.init(extents, ghosts, hasBC),
2045  operand5_.init(extents, ghosts, hasBC),
2046  operand6_.init(extents, ghosts, hasBC));
2047  }
2048 
2049  private:
2050  FunctorFactory const & functorFactory_;
2051 
2052  Operand1 const operand1_;
2053 
2054  Operand2 const operand2_;
2055 
2056  Operand3 const operand3_;
2057 
2058  Operand4 const operand4_;
2059 
2060  Operand5 const operand5_;
2061 
2062  Operand6 const operand6_;
2063  }
2064  #endif
2065  /* ENABLE_THREADS */;
2066  template<typename FunctorFactory,
2067  typename Operand1,
2068  typename Operand2,
2069  typename Operand3,
2070  typename Operand4,
2071  typename Operand5,
2072  typename Operand6>
2073  struct ApplyPointwise6<SeqWalk,
2074  FunctorFactory,
2075  Operand1,
2076  Operand2,
2077  Operand3,
2078  Operand4,
2079  Operand5,
2080  Operand6> {
2081  public:
2082  typename Operand1::value_type typedef value_type;
2083 
2084  ApplyPointwise6(FunctorFactory const & functorFactory,
2085  Operand1 const & operand1,
2086  Operand2 const & operand2,
2087  Operand3 const & operand3,
2088  Operand4 const & operand4,
2089  Operand5 const & operand5,
2090  Operand6 const & operand6)
2091  : functor_(functorFactory()),
2092  operand1_(operand1),
2093  operand2_(operand2),
2094  operand3_(operand3),
2095  operand4_(operand4),
2096  operand5_(operand5),
2097  operand6_(operand6)
2098  {}
2099 
2100  inline value_type eval(int const x, int const y, int const z) const {
2101  return functor_(operand1_.eval(x, y, z),
2102  operand2_.eval(x, y, z),
2103  operand3_.eval(x, y, z),
2104  operand4_.eval(x, y, z),
2105  operand5_.eval(x, y, z),
2106  operand6_.eval(x, y, z));
2107  }
2108 
2109  private:
2110  typename FunctorFactory::Operator functor_;
2111 
2112  Operand1 operand1_;
2113 
2114  Operand2 operand2_;
2115 
2116  Operand3 operand3_;
2117 
2118  Operand4 operand4_;
2119 
2120  Operand5 operand5_;
2121 
2122  Operand6 operand6_;
2123  };
2124  #ifdef __CUDACC__
2125  template<typename FunctorFactory,
2126  typename Operand1,
2127  typename Operand2,
2128  typename Operand3,
2129  typename Operand4,
2130  typename Operand5,
2131  typename Operand6>
2132  struct ApplyPointwise6<GPUWalk,
2133  FunctorFactory,
2134  Operand1,
2135  Operand2,
2136  Operand3,
2137  Operand4,
2138  Operand5,
2139  Operand6> {
2140  public:
2141  typename Operand1::value_type typedef value_type;
2142 
2143  ApplyPointwise6(Operand1 const & operand1,
2144  Operand2 const & operand2,
2145  Operand3 const & operand3,
2146  Operand4 const & operand4,
2147  Operand5 const & operand5,
2148  Operand6 const & operand6)
2149  : operand1_(operand1),
2150  operand2_(operand2),
2151  operand3_(operand3),
2152  operand4_(operand4),
2153  operand5_(operand5),
2154  operand6_(operand6)
2155  {}
2156 
2157  __device__ inline value_type eval(int const x,
2158  int const y,
2159  int const z) const { return 0; }
2160 
2161  private:
2162  Operand1 operand1_;
2163 
2164  Operand2 operand2_;
2165 
2166  Operand3 operand3_;
2167 
2168  Operand4 operand4_;
2169 
2170  Operand5 operand5_;
2171 
2172  Operand6 operand6_;
2173  }
2174  #endif
2175  /* __CUDACC__ */;
2176 
2177  template<typename Operand1,
2178  typename Operand2,
2179  typename Operand3,
2180  typename Operand4,
2181  typename Operand5,
2182  typename Operand6>
2185  typename RefineFieldType<typename
2187  Result,
2188  typename
2189  RefineFieldType<typename
2191  Result,
2192  typename
2193  RefineFieldType<typename
2195  Result,
2196  typename
2197  RefineFieldType<typename
2199  Result,
2200  typename
2202  Result>::
2203  Result>::
2204  Result>::
2205  Result>::Result>::
2206  Result typedef Result;
2207  };
2208 
2209  template<typename FunctorFactory,
2210  typename Operand1,
2211  typename Operand2,
2212  typename Operand3,
2213  typename Operand4,
2214  typename Operand5,
2215  typename Operand6>
2216  inline NeboExpression<ApplyPointwise6<Initial,
2217  FunctorFactory,
2218  typename FinalType<Operand1>::
2219  Result,
2220  typename FinalType<Operand2>::
2221  Result,
2222  typename FinalType<Operand3>::
2223  Result,
2224  typename FinalType<Operand4>::
2225  Result,
2226  typename FinalType<Operand5>::
2227  Result,
2228  typename FinalType<Operand6>::
2229  Result>,
2230  typename FindFieldType6<Operand1,
2231  Operand2,
2232  Operand3,
2233  Operand4,
2234  Operand5,
2235  Operand6>::Result>
2236  apply_pointwise(FunctorFactory const & functorFactory,
2237  Operand1 operand1,
2238  Operand2 operand2,
2239  Operand3 operand3,
2240  Operand4 operand4,
2241  Operand5 operand5,
2242  Operand6 operand6) {
2243  ApplyPointwise6<Initial,
2244  FunctorFactory,
2245  typename FinalType<Operand1>::Result,
2246  typename FinalType<Operand2>::Result,
2247  typename FinalType<Operand3>::Result,
2248  typename FinalType<Operand4>::Result,
2249  typename FinalType<Operand5>::Result,
2250  typename FinalType<Operand6>::Result> typedef
2251  ReturnType;
2252 
2253  NeboExpression<ReturnType,
2254  typename FindFieldType6<Operand1,
2255  Operand2,
2256  Operand3,
2257  Operand4,
2258  Operand5,
2259  Operand6>::Result> typedef
2260  ReturnTerm;
2261 
2262  return ReturnTerm(ReturnType(functorFactory,
2263  normalize(operand1),
2264  normalize(operand2),
2265  normalize(operand3),
2266  normalize(operand4),
2267  normalize(operand5),
2268  normalize(operand6)));
2269  }
2270 
2271  template<typename CurrentMode,
2272  typename FunctorFactory,
2273  typename Operand1,
2274  typename Operand2,
2275  typename Operand3,
2276  typename Operand4,
2277  typename Operand5,
2278  typename Operand6,
2279  typename Operand7>
2281  template<typename FunctorFactory,
2282  typename Operand1,
2283  typename Operand2,
2284  typename Operand3,
2285  typename Operand4,
2286  typename Operand5,
2287  typename Operand6,
2288  typename Operand7>
2289  struct ApplyPointwise7<Initial,
2290  FunctorFactory,
2291  Operand1,
2292  Operand2,
2293  Operand3,
2294  Operand4,
2295  Operand5,
2296  Operand6,
2297  Operand7> {
2298  public:
2299  ApplyPointwise7<SeqWalk,
2300  FunctorFactory,
2301  typename Operand1::SeqWalkType,
2302  typename Operand2::SeqWalkType,
2303  typename Operand3::SeqWalkType,
2304  typename Operand4::SeqWalkType,
2305  typename Operand5::SeqWalkType,
2306  typename Operand6::SeqWalkType,
2307  typename Operand7::SeqWalkType> typedef SeqWalkType;
2308 
2309  #ifdef ENABLE_THREADS
2310  ApplyPointwise7<Resize,
2311  FunctorFactory,
2312  typename Operand1::ResizeType,
2313  typename Operand2::ResizeType,
2314  typename Operand3::ResizeType,
2315  typename Operand4::ResizeType,
2316  typename Operand5::ResizeType,
2317  typename Operand6::ResizeType,
2318  typename Operand7::ResizeType> typedef ResizeType;
2319  #endif
2320  /* ENABLE_THREADS */
2321 
2322  #ifdef __CUDACC__
2323  ApplyPointwise7<GPUWalk,
2324  FunctorFactory,
2325  typename Operand1::GPUWalkType,
2326  typename Operand2::GPUWalkType,
2327  typename Operand3::GPUWalkType,
2328  typename Operand4::GPUWalkType,
2329  typename Operand5::GPUWalkType,
2330  typename Operand6::GPUWalkType,
2331  typename Operand7::GPUWalkType> typedef GPUWalkType
2332  ;
2333  #endif
2334  /* __CUDACC__ */
2335 
2336  ApplyPointwise7(FunctorFactory const & functorFactory,
2337  Operand1 const & operand1,
2338  Operand2 const & operand2,
2339  Operand3 const & operand3,
2340  Operand4 const & operand4,
2341  Operand5 const & operand5,
2342  Operand6 const & operand6,
2343  Operand7 const & operand7)
2344  : functorFactory_(functorFactory),
2345  operand1_(operand1),
2346  operand2_(operand2),
2347  operand3_(operand3),
2348  operand4_(operand4),
2349  operand5_(operand5),
2350  operand6_(operand6),
2351  operand7_(operand7)
2352  {}
2353 
2354  inline GhostData ghosts_with_bc(void) const {
2355  return min(operand1_.ghosts_with_bc(),
2356  min(operand2_.ghosts_with_bc(),
2357  min(operand3_.ghosts_with_bc(),
2358  min(operand4_.ghosts_with_bc(),
2359  min(operand5_.ghosts_with_bc(),
2360  min(operand6_.ghosts_with_bc(),
2361  operand7_.ghosts_with_bc()))))));
2362  }
2363 
2364  inline GhostData ghosts_without_bc(void) const {
2365  return min(operand1_.ghosts_without_bc(),
2366  min(operand2_.ghosts_without_bc(),
2367  min(operand3_.ghosts_without_bc(),
2368  min(operand4_.ghosts_without_bc(),
2369  min(operand5_.ghosts_without_bc(),
2370  min(operand6_.ghosts_without_bc(),
2371  operand7_.ghosts_without_bc()))))));
2372  }
2373 
2374  inline bool has_extents(void) const {
2375  return (operand1_.has_extents() || operand2_.has_extents() ||
2376  operand3_.has_extents() || operand4_.has_extents() ||
2377  operand5_.has_extents() || operand6_.has_extents() ||
2378  operand7_.has_extents());
2379  }
2380 
2381  inline IntVec extents(void) const {
2382  #ifndef NDEBUG
2383  if((operand1_.has_extents() || operand2_.has_extents() ||
2384  operand3_.has_extents() || operand4_.has_extents() ||
2385  operand5_.has_extents() || operand6_.has_extents() ||
2386  operand7_.has_extents())) {
2387  IntVec extents;
2388 
2389  if(operand1_.has_extents()) { extents = operand1_.extents(); }
2390  else {
2391  if(operand2_.has_extents()) {
2392  extents = operand2_.extents();
2393  }
2394  else {
2395  if(operand3_.has_extents()) {
2396  extents = operand3_.extents();
2397  }
2398  else {
2399  if(operand4_.has_extents()) {
2400  extents = operand4_.extents();
2401  }
2402  else {
2403  if(operand5_.has_extents()) {
2404  extents = operand5_.extents();
2405  }
2406  else {
2407  if(operand6_.has_extents()) {
2408  extents = operand6_.extents();
2409  }
2410  else { extents = operand7_.extents(); };
2411  };
2412  };
2413  };
2414  };
2415  };
2416 
2417  if(operand1_.has_extents()) {
2418  assert(extents == operand1_.extents());
2419  };
2420 
2421  if(operand2_.has_extents()) {
2422  assert(extents == operand2_.extents());
2423  };
2424 
2425  if(operand3_.has_extents()) {
2426  assert(extents == operand3_.extents());
2427  };
2428 
2429  if(operand4_.has_extents()) {
2430  assert(extents == operand4_.extents());
2431  };
2432 
2433  if(operand5_.has_extents()) {
2434  assert(extents == operand5_.extents());
2435  };
2436 
2437  if(operand6_.has_extents()) {
2438  assert(extents == operand6_.extents());
2439  };
2440 
2441  if(operand7_.has_extents()) {
2442  assert(extents == operand7_.extents());
2443  };
2444  }
2445  #endif
2446  /* NDEBUG */;
2447 
2448  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
2449  ?
2450  operand2_.extents()
2451  : ((operand3_.has_extents()
2452  ?
2453  operand3_.extents()
2454  : ((operand4_.has_extents()
2455  ?
2456  operand4_.extents()
2457  :
2458  ((operand5_.has_extents()
2459  ?
2460  operand5_.extents()
2461  :
2462  ((operand6_.has_extents()
2463  ?
2464  operand6_.extents()
2465  :
2466  (operand7_.extents()))))))))))));
2467  }
2468 
2469  inline IntVec has_bc(void) const {
2470  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
2471  ?
2472  operand2_.has_bc()
2473  : ((operand3_.has_extents()
2474  ?
2475  operand3_.has_bc()
2476  : ((operand4_.has_extents()
2477  ?
2478  operand4_.has_bc()
2479  : ((operand5_.has_extents()
2480  ?
2481  operand5_.has_bc()
2482  :
2483  ((operand6_.has_extents()
2484  ?
2485  operand6_.has_bc()
2486  :
2487  (operand7_.has_bc()))))))))))));
2488  }
2489 
2490  inline SeqWalkType init(IntVec const & extents,
2491  GhostData const & ghosts,
2492  IntVec const & hasBC) const {
2493  return SeqWalkType(functorFactory_,
2494  operand1_.init(extents, ghosts, hasBC),
2495  operand2_.init(extents, ghosts, hasBC),
2496  operand3_.init(extents, ghosts, hasBC),
2497  operand4_.init(extents, ghosts, hasBC),
2498  operand5_.init(extents, ghosts, hasBC),
2499  operand6_.init(extents, ghosts, hasBC),
2500  operand7_.init(extents, ghosts, hasBC));
2501  }
2502 
2503  #ifdef ENABLE_THREADS
2504  inline ResizeType resize(void) const {
2505  return ResizeType(functorFactory_,
2506  operand1_.resize(),
2507  operand2_.resize(),
2508  operand3_.resize(),
2509  operand4_.resize(),
2510  operand5_.resize(),
2511  operand6_.resize(),
2512  operand7_.resize());
2513  }
2514  #endif
2515  /* ENABLE_THREADS */
2516 
2517  #ifdef __CUDACC__
2518  inline bool cpu_ready(void) const {
2519  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
2520  operand3_.cpu_ready() && operand4_.cpu_ready() &&
2521  operand5_.cpu_ready() && operand6_.cpu_ready() &&
2522  operand7_.cpu_ready());
2523  }
2524 
2525  inline bool gpu_ready(int const deviceIndex) const { return false; }
2526 
2527  inline GPUWalkType gpu_init(IntVec const & extents,
2528  GhostData const & ghosts,
2529  IntVec const & hasBC,
2530  int const deviceIndex,
2531  cudaStream_t const & lhsStream) const {
2532  return GPUWalkType(operand1_.gpu_init(extents,
2533  ghosts,
2534  hasBC,
2535  deviceIndex,
2536  lhsStream),
2537  operand2_.gpu_init(extents,
2538  ghosts,
2539  hasBC,
2540  deviceIndex,
2541  lhsStream),
2542  operand3_.gpu_init(extents,
2543  ghosts,
2544  hasBC,
2545  deviceIndex,
2546  lhsStream),
2547  operand4_.gpu_init(extents,
2548  ghosts,
2549  hasBC,
2550  deviceIndex,
2551  lhsStream),
2552  operand5_.gpu_init(extents,
2553  ghosts,
2554  hasBC,
2555  deviceIndex,
2556  lhsStream),
2557  operand6_.gpu_init(extents,
2558  ghosts,
2559  hasBC,
2560  deviceIndex,
2561  lhsStream),
2562  operand7_.gpu_init(extents,
2563  ghosts,
2564  hasBC,
2565  deviceIndex,
2566  lhsStream));
2567  }
2568 
2569  inline void stream_wait_event(cudaEvent_t const & event) const {
2570  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
2571  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
2572  operand5_.stream_wait_event(event); operand6_.stream_wait_event(event);
2573  operand7_.stream_wait_event(event);
2574  }
2575 
2576  #ifdef NEBO_GPU_TEST
2577  inline void gpu_prep(int const deviceIndex) const {
2578  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
2579  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
2580  operand5_.gpu_prep(deviceIndex); operand6_.gpu_prep(deviceIndex);
2581  operand7_.gpu_prep(deviceIndex);
2582  }
2583  #endif
2584  /* NEBO_GPU_TEST */
2585  #endif
2586  /* __CUDACC__ */
2587 
2588  private:
2589  FunctorFactory const & functorFactory_;
2590 
2591  Operand1 const operand1_;
2592 
2593  Operand2 const operand2_;
2594 
2595  Operand3 const operand3_;
2596 
2597  Operand4 const operand4_;
2598 
2599  Operand5 const operand5_;
2600 
2601  Operand6 const operand6_;
2602 
2603  Operand7 const operand7_;
2604  };
2605  #ifdef ENABLE_THREADS
2606  template<typename FunctorFactory,
2607  typename Operand1,
2608  typename Operand2,
2609  typename Operand3,
2610  typename Operand4,
2611  typename Operand5,
2612  typename Operand6,
2613  typename Operand7>
2614  struct ApplyPointwise7<Resize,
2615  FunctorFactory,
2616  Operand1,
2617  Operand2,
2618  Operand3,
2619  Operand4,
2620  Operand5,
2621  Operand6,
2622  Operand7> {
2623  public:
2624  ApplyPointwise7<SeqWalk,
2625  FunctorFactory,
2626  typename Operand1::SeqWalkType,
2627  typename Operand2::SeqWalkType,
2628  typename Operand3::SeqWalkType,
2629  typename Operand4::SeqWalkType,
2630  typename Operand5::SeqWalkType,
2631  typename Operand6::SeqWalkType,
2632  typename Operand7::SeqWalkType> typedef SeqWalkType
2633  ;
2634 
2635  ApplyPointwise7(FunctorFactory const & functorFactory,
2636  Operand1 const & operand1,
2637  Operand2 const & operand2,
2638  Operand3 const & operand3,
2639  Operand4 const & operand4,
2640  Operand5 const & operand5,
2641  Operand6 const & operand6,
2642  Operand7 const & operand7)
2643  : functorFactory_(functorFactory),
2644  operand1_(operand1),
2645  operand2_(operand2),
2646  operand3_(operand3),
2647  operand4_(operand4),
2648  operand5_(operand5),
2649  operand6_(operand6),
2650  operand7_(operand7)
2651  {}
2652 
2653  inline SeqWalkType init(IntVec const & extents,
2654  GhostData const & ghosts,
2655  IntVec const & hasBC) const {
2656  return SeqWalkType(functorFactory_,
2657  operand1_.init(extents, ghosts, hasBC),
2658  operand2_.init(extents, ghosts, hasBC),
2659  operand3_.init(extents, ghosts, hasBC),
2660  operand4_.init(extents, ghosts, hasBC),
2661  operand5_.init(extents, ghosts, hasBC),
2662  operand6_.init(extents, ghosts, hasBC),
2663  operand7_.init(extents, ghosts, hasBC));
2664  }
2665 
2666  private:
2667  FunctorFactory const & functorFactory_;
2668 
2669  Operand1 const operand1_;
2670 
2671  Operand2 const operand2_;
2672 
2673  Operand3 const operand3_;
2674 
2675  Operand4 const operand4_;
2676 
2677  Operand5 const operand5_;
2678 
2679  Operand6 const operand6_;
2680 
2681  Operand7 const operand7_;
2682  }
2683  #endif
2684  /* ENABLE_THREADS */;
2685  template<typename FunctorFactory,
2686  typename Operand1,
2687  typename Operand2,
2688  typename Operand3,
2689  typename Operand4,
2690  typename Operand5,
2691  typename Operand6,
2692  typename Operand7>
2693  struct ApplyPointwise7<SeqWalk,
2694  FunctorFactory,
2695  Operand1,
2696  Operand2,
2697  Operand3,
2698  Operand4,
2699  Operand5,
2700  Operand6,
2701  Operand7> {
2702  public:
2703  typename Operand1::value_type typedef value_type;
2704 
2705  ApplyPointwise7(FunctorFactory const & functorFactory,
2706  Operand1 const & operand1,
2707  Operand2 const & operand2,
2708  Operand3 const & operand3,
2709  Operand4 const & operand4,
2710  Operand5 const & operand5,
2711  Operand6 const & operand6,
2712  Operand7 const & operand7)
2713  : functor_(functorFactory()),
2714  operand1_(operand1),
2715  operand2_(operand2),
2716  operand3_(operand3),
2717  operand4_(operand4),
2718  operand5_(operand5),
2719  operand6_(operand6),
2720  operand7_(operand7)
2721  {}
2722 
2723  inline value_type eval(int const x, int const y, int const z) const {
2724  return functor_(operand1_.eval(x, y, z),
2725  operand2_.eval(x, y, z),
2726  operand3_.eval(x, y, z),
2727  operand4_.eval(x, y, z),
2728  operand5_.eval(x, y, z),
2729  operand6_.eval(x, y, z),
2730  operand7_.eval(x, y, z));
2731  }
2732 
2733  private:
2734  typename FunctorFactory::Operator functor_;
2735 
2736  Operand1 operand1_;
2737 
2738  Operand2 operand2_;
2739 
2740  Operand3 operand3_;
2741 
2742  Operand4 operand4_;
2743 
2744  Operand5 operand5_;
2745 
2746  Operand6 operand6_;
2747 
2748  Operand7 operand7_;
2749  };
2750  #ifdef __CUDACC__
2751  template<typename FunctorFactory,
2752  typename Operand1,
2753  typename Operand2,
2754  typename Operand3,
2755  typename Operand4,
2756  typename Operand5,
2757  typename Operand6,
2758  typename Operand7>
2759  struct ApplyPointwise7<GPUWalk,
2760  FunctorFactory,
2761  Operand1,
2762  Operand2,
2763  Operand3,
2764  Operand4,
2765  Operand5,
2766  Operand6,
2767  Operand7> {
2768  public:
2769  typename Operand1::value_type typedef value_type;
2770 
2771  ApplyPointwise7(Operand1 const & operand1,
2772  Operand2 const & operand2,
2773  Operand3 const & operand3,
2774  Operand4 const & operand4,
2775  Operand5 const & operand5,
2776  Operand6 const & operand6,
2777  Operand7 const & operand7)
2778  : operand1_(operand1),
2779  operand2_(operand2),
2780  operand3_(operand3),
2781  operand4_(operand4),
2782  operand5_(operand5),
2783  operand6_(operand6),
2784  operand7_(operand7)
2785  {}
2786 
2787  __device__ inline value_type eval(int const x,
2788  int const y,
2789  int const z) const { return 0; }
2790 
2791  private:
2792  Operand1 operand1_;
2793 
2794  Operand2 operand2_;
2795 
2796  Operand3 operand3_;
2797 
2798  Operand4 operand4_;
2799 
2800  Operand5 operand5_;
2801 
2802  Operand6 operand6_;
2803 
2804  Operand7 operand7_;
2805  }
2806  #endif
2807  /* __CUDACC__ */;
2808 
2809  template<typename Operand1,
2810  typename Operand2,
2811  typename Operand3,
2812  typename Operand4,
2813  typename Operand5,
2814  typename Operand6,
2815  typename Operand7>
2818  typename RefineFieldType<typename
2820  Result,
2821  typename
2822  RefineFieldType<typename
2824  Result,
2825  typename
2826  RefineFieldType<typename
2828  Result,
2829  typename
2830  RefineFieldType<typename
2832  Result,
2833  typename
2834  RefineFieldType<typename
2836  Result,
2837  typename
2839  Result>::
2840  Result>::
2841  Result>::
2842  Result>::
2843  Result>::Result>::
2844  Result typedef Result;
2845  };
2846 
2847  template<typename FunctorFactory,
2848  typename Operand1,
2849  typename Operand2,
2850  typename Operand3,
2851  typename Operand4,
2852  typename Operand5,
2853  typename Operand6,
2854  typename Operand7>
2855  inline NeboExpression<ApplyPointwise7<Initial,
2856  FunctorFactory,
2857  typename FinalType<Operand1>::
2858  Result,
2859  typename FinalType<Operand2>::
2860  Result,
2861  typename FinalType<Operand3>::
2862  Result,
2863  typename FinalType<Operand4>::
2864  Result,
2865  typename FinalType<Operand5>::
2866  Result,
2867  typename FinalType<Operand6>::
2868  Result,
2869  typename FinalType<Operand7>::
2870  Result>,
2871  typename FindFieldType7<Operand1,
2872  Operand2,
2873  Operand3,
2874  Operand4,
2875  Operand5,
2876  Operand6,
2877  Operand7>::Result>
2878  apply_pointwise(FunctorFactory const & functorFactory,
2879  Operand1 operand1,
2880  Operand2 operand2,
2881  Operand3 operand3,
2882  Operand4 operand4,
2883  Operand5 operand5,
2884  Operand6 operand6,
2885  Operand7 operand7) {
2886  ApplyPointwise7<Initial,
2887  FunctorFactory,
2888  typename FinalType<Operand1>::Result,
2889  typename FinalType<Operand2>::Result,
2890  typename FinalType<Operand3>::Result,
2891  typename FinalType<Operand4>::Result,
2892  typename FinalType<Operand5>::Result,
2893  typename FinalType<Operand6>::Result,
2894  typename FinalType<Operand7>::Result> typedef
2895  ReturnType;
2896 
2897  NeboExpression<ReturnType,
2898  typename FindFieldType7<Operand1,
2899  Operand2,
2900  Operand3,
2901  Operand4,
2902  Operand5,
2903  Operand6,
2904  Operand7>::Result> typedef
2905  ReturnTerm;
2906 
2907  return ReturnTerm(ReturnType(functorFactory,
2908  normalize(operand1),
2909  normalize(operand2),
2910  normalize(operand3),
2911  normalize(operand4),
2912  normalize(operand5),
2913  normalize(operand6),
2914  normalize(operand7)));
2915  }
2916 
2917  template<typename CurrentMode,
2918  typename FunctorFactory,
2919  typename Operand1,
2920  typename Operand2,
2921  typename Operand3,
2922  typename Operand4,
2923  typename Operand5,
2924  typename Operand6,
2925  typename Operand7,
2926  typename Operand8>
2928  template<typename FunctorFactory,
2929  typename Operand1,
2930  typename Operand2,
2931  typename Operand3,
2932  typename Operand4,
2933  typename Operand5,
2934  typename Operand6,
2935  typename Operand7,
2936  typename Operand8>
2937  struct ApplyPointwise8<Initial,
2938  FunctorFactory,
2939  Operand1,
2940  Operand2,
2941  Operand3,
2942  Operand4,
2943  Operand5,
2944  Operand6,
2945  Operand7,
2946  Operand8> {
2947  public:
2948  ApplyPointwise8<SeqWalk,
2949  FunctorFactory,
2950  typename Operand1::SeqWalkType,
2951  typename Operand2::SeqWalkType,
2952  typename Operand3::SeqWalkType,
2953  typename Operand4::SeqWalkType,
2954  typename Operand5::SeqWalkType,
2955  typename Operand6::SeqWalkType,
2956  typename Operand7::SeqWalkType,
2957  typename Operand8::SeqWalkType> typedef SeqWalkType;
2958 
2959  #ifdef ENABLE_THREADS
2960  ApplyPointwise8<Resize,
2961  FunctorFactory,
2962  typename Operand1::ResizeType,
2963  typename Operand2::ResizeType,
2964  typename Operand3::ResizeType,
2965  typename Operand4::ResizeType,
2966  typename Operand5::ResizeType,
2967  typename Operand6::ResizeType,
2968  typename Operand7::ResizeType,
2969  typename Operand8::ResizeType> typedef ResizeType;
2970  #endif
2971  /* ENABLE_THREADS */
2972 
2973  #ifdef __CUDACC__
2974  ApplyPointwise8<GPUWalk,
2975  FunctorFactory,
2976  typename Operand1::GPUWalkType,
2977  typename Operand2::GPUWalkType,
2978  typename Operand3::GPUWalkType,
2979  typename Operand4::GPUWalkType,
2980  typename Operand5::GPUWalkType,
2981  typename Operand6::GPUWalkType,
2982  typename Operand7::GPUWalkType,
2983  typename Operand8::GPUWalkType> typedef GPUWalkType
2984  ;
2985  #endif
2986  /* __CUDACC__ */
2987 
2988  ApplyPointwise8(FunctorFactory const & functorFactory,
2989  Operand1 const & operand1,
2990  Operand2 const & operand2,
2991  Operand3 const & operand3,
2992  Operand4 const & operand4,
2993  Operand5 const & operand5,
2994  Operand6 const & operand6,
2995  Operand7 const & operand7,
2996  Operand8 const & operand8)
2997  : functorFactory_(functorFactory),
2998  operand1_(operand1),
2999  operand2_(operand2),
3000  operand3_(operand3),
3001  operand4_(operand4),
3002  operand5_(operand5),
3003  operand6_(operand6),
3004  operand7_(operand7),
3005  operand8_(operand8)
3006  {}
3007 
3008  inline GhostData ghosts_with_bc(void) const {
3009  return min(operand1_.ghosts_with_bc(),
3010  min(operand2_.ghosts_with_bc(),
3011  min(operand3_.ghosts_with_bc(),
3012  min(operand4_.ghosts_with_bc(),
3013  min(operand5_.ghosts_with_bc(),
3014  min(operand6_.ghosts_with_bc(),
3015  min(operand7_.ghosts_with_bc(),
3016  operand8_.ghosts_with_bc())))))));
3017  }
3018 
3019  inline GhostData ghosts_without_bc(void) const {
3020  return min(operand1_.ghosts_without_bc(),
3021  min(operand2_.ghosts_without_bc(),
3022  min(operand3_.ghosts_without_bc(),
3023  min(operand4_.ghosts_without_bc(),
3024  min(operand5_.ghosts_without_bc(),
3025  min(operand6_.ghosts_without_bc(),
3026  min(operand7_.ghosts_without_bc(),
3027  operand8_.ghosts_without_bc())))))));
3028  }
3029 
3030  inline bool has_extents(void) const {
3031  return (operand1_.has_extents() || operand2_.has_extents() ||
3032  operand3_.has_extents() || operand4_.has_extents() ||
3033  operand5_.has_extents() || operand6_.has_extents() ||
3034  operand7_.has_extents() || operand8_.has_extents());
3035  }
3036 
3037  inline IntVec extents(void) const {
3038  #ifndef NDEBUG
3039  if((operand1_.has_extents() || operand2_.has_extents() ||
3040  operand3_.has_extents() || operand4_.has_extents() ||
3041  operand5_.has_extents() || operand6_.has_extents() ||
3042  operand7_.has_extents() || operand8_.has_extents())) {
3043  IntVec extents;
3044 
3045  if(operand1_.has_extents()) { extents = operand1_.extents(); }
3046  else {
3047  if(operand2_.has_extents()) {
3048  extents = operand2_.extents();
3049  }
3050  else {
3051  if(operand3_.has_extents()) {
3052  extents = operand3_.extents();
3053  }
3054  else {
3055  if(operand4_.has_extents()) {
3056  extents = operand4_.extents();
3057  }
3058  else {
3059  if(operand5_.has_extents()) {
3060  extents = operand5_.extents();
3061  }
3062  else {
3063  if(operand6_.has_extents()) {
3064  extents = operand6_.extents();
3065  }
3066  else {
3067  if(operand7_.has_extents()) {
3068  extents = operand7_.extents();
3069  }
3070  else { extents = operand8_.extents(); };
3071  };
3072  };
3073  };
3074  };
3075  };
3076  };
3077 
3078  if(operand1_.has_extents()) {
3079  assert(extents == operand1_.extents());
3080  };
3081 
3082  if(operand2_.has_extents()) {
3083  assert(extents == operand2_.extents());
3084  };
3085 
3086  if(operand3_.has_extents()) {
3087  assert(extents == operand3_.extents());
3088  };
3089 
3090  if(operand4_.has_extents()) {
3091  assert(extents == operand4_.extents());
3092  };
3093 
3094  if(operand5_.has_extents()) {
3095  assert(extents == operand5_.extents());
3096  };
3097 
3098  if(operand6_.has_extents()) {
3099  assert(extents == operand6_.extents());
3100  };
3101 
3102  if(operand7_.has_extents()) {
3103  assert(extents == operand7_.extents());
3104  };
3105 
3106  if(operand8_.has_extents()) {
3107  assert(extents == operand8_.extents());
3108  };
3109  }
3110  #endif
3111  /* NDEBUG */;
3112 
3113  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
3114  ?
3115  operand2_.extents()
3116  : ((operand3_.has_extents()
3117  ?
3118  operand3_.extents()
3119  : ((operand4_.has_extents()
3120  ?
3121  operand4_.extents()
3122  :
3123  ((operand5_.has_extents()
3124  ?
3125  operand5_.extents()
3126  :
3127  ((operand6_.has_extents()
3128  ?
3129  operand6_.extents()
3130  :
3131  ((operand7_.has_extents()
3132  ?
3133  operand7_.extents()
3134  :
3135  (operand8_.extents()))))))))))))));
3136  }
3137 
3138  inline IntVec has_bc(void) const {
3139  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
3140  ?
3141  operand2_.has_bc()
3142  : ((operand3_.has_extents()
3143  ?
3144  operand3_.has_bc()
3145  : ((operand4_.has_extents()
3146  ?
3147  operand4_.has_bc()
3148  : ((operand5_.has_extents()
3149  ?
3150  operand5_.has_bc()
3151  :
3152  ((operand6_.has_extents()
3153  ?
3154  operand6_.has_bc()
3155  :
3156  ((operand7_.has_extents()
3157  ?
3158  operand7_.has_bc()
3159  :
3160  (operand8_.has_bc()))))))))))))));
3161  }
3162 
3163  inline SeqWalkType init(IntVec const & extents,
3164  GhostData const & ghosts,
3165  IntVec const & hasBC) const {
3166  return SeqWalkType(functorFactory_,
3167  operand1_.init(extents, ghosts, hasBC),
3168  operand2_.init(extents, ghosts, hasBC),
3169  operand3_.init(extents, ghosts, hasBC),
3170  operand4_.init(extents, ghosts, hasBC),
3171  operand5_.init(extents, ghosts, hasBC),
3172  operand6_.init(extents, ghosts, hasBC),
3173  operand7_.init(extents, ghosts, hasBC),
3174  operand8_.init(extents, ghosts, hasBC));
3175  }
3176 
3177  #ifdef ENABLE_THREADS
3178  inline ResizeType resize(void) const {
3179  return ResizeType(functorFactory_,
3180  operand1_.resize(),
3181  operand2_.resize(),
3182  operand3_.resize(),
3183  operand4_.resize(),
3184  operand5_.resize(),
3185  operand6_.resize(),
3186  operand7_.resize(),
3187  operand8_.resize());
3188  }
3189  #endif
3190  /* ENABLE_THREADS */
3191 
3192  #ifdef __CUDACC__
3193  inline bool cpu_ready(void) const {
3194  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
3195  operand3_.cpu_ready() && operand4_.cpu_ready() &&
3196  operand5_.cpu_ready() && operand6_.cpu_ready() &&
3197  operand7_.cpu_ready() && operand8_.cpu_ready());
3198  }
3199 
3200  inline bool gpu_ready(int const deviceIndex) const { return false; }
3201 
3202  inline GPUWalkType gpu_init(IntVec const & extents,
3203  GhostData const & ghosts,
3204  IntVec const & hasBC,
3205  int const deviceIndex,
3206  cudaStream_t const & lhsStream) const {
3207  return GPUWalkType(operand1_.gpu_init(extents,
3208  ghosts,
3209  hasBC,
3210  deviceIndex,
3211  lhsStream),
3212  operand2_.gpu_init(extents,
3213  ghosts,
3214  hasBC,
3215  deviceIndex,
3216  lhsStream),
3217  operand3_.gpu_init(extents,
3218  ghosts,
3219  hasBC,
3220  deviceIndex,
3221  lhsStream),
3222  operand4_.gpu_init(extents,
3223  ghosts,
3224  hasBC,
3225  deviceIndex,
3226  lhsStream),
3227  operand5_.gpu_init(extents,
3228  ghosts,
3229  hasBC,
3230  deviceIndex,
3231  lhsStream),
3232  operand6_.gpu_init(extents,
3233  ghosts,
3234  hasBC,
3235  deviceIndex,
3236  lhsStream),
3237  operand7_.gpu_init(extents,
3238  ghosts,
3239  hasBC,
3240  deviceIndex,
3241  lhsStream),
3242  operand8_.gpu_init(extents,
3243  ghosts,
3244  hasBC,
3245  deviceIndex,
3246  lhsStream));
3247  }
3248 
3249  inline void stream_wait_event(cudaEvent_t const & event) const {
3250  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
3251  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
3252  operand5_.stream_wait_event(event); operand6_.stream_wait_event(event);
3253  operand7_.stream_wait_event(event); operand8_.stream_wait_event(event);
3254  }
3255 
3256  #ifdef NEBO_GPU_TEST
3257  inline void gpu_prep(int const deviceIndex) const {
3258  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
3259  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
3260  operand5_.gpu_prep(deviceIndex); operand6_.gpu_prep(deviceIndex);
3261  operand7_.gpu_prep(deviceIndex); operand8_.gpu_prep(deviceIndex);
3262  }
3263  #endif
3264  /* NEBO_GPU_TEST */
3265  #endif
3266  /* __CUDACC__ */
3267 
3268  private:
3269  FunctorFactory const & functorFactory_;
3270 
3271  Operand1 const operand1_;
3272 
3273  Operand2 const operand2_;
3274 
3275  Operand3 const operand3_;
3276 
3277  Operand4 const operand4_;
3278 
3279  Operand5 const operand5_;
3280 
3281  Operand6 const operand6_;
3282 
3283  Operand7 const operand7_;
3284 
3285  Operand8 const operand8_;
3286  };
3287  #ifdef ENABLE_THREADS
3288  template<typename FunctorFactory,
3289  typename Operand1,
3290  typename Operand2,
3291  typename Operand3,
3292  typename Operand4,
3293  typename Operand5,
3294  typename Operand6,
3295  typename Operand7,
3296  typename Operand8>
3297  struct ApplyPointwise8<Resize,
3298  FunctorFactory,
3299  Operand1,
3300  Operand2,
3301  Operand3,
3302  Operand4,
3303  Operand5,
3304  Operand6,
3305  Operand7,
3306  Operand8> {
3307  public:
3308  ApplyPointwise8<SeqWalk,
3309  FunctorFactory,
3310  typename Operand1::SeqWalkType,
3311  typename Operand2::SeqWalkType,
3312  typename Operand3::SeqWalkType,
3313  typename Operand4::SeqWalkType,
3314  typename Operand5::SeqWalkType,
3315  typename Operand6::SeqWalkType,
3316  typename Operand7::SeqWalkType,
3317  typename Operand8::SeqWalkType> typedef SeqWalkType
3318  ;
3319 
3320  ApplyPointwise8(FunctorFactory const & functorFactory,
3321  Operand1 const & operand1,
3322  Operand2 const & operand2,
3323  Operand3 const & operand3,
3324  Operand4 const & operand4,
3325  Operand5 const & operand5,
3326  Operand6 const & operand6,
3327  Operand7 const & operand7,
3328  Operand8 const & operand8)
3329  : functorFactory_(functorFactory),
3330  operand1_(operand1),
3331  operand2_(operand2),
3332  operand3_(operand3),
3333  operand4_(operand4),
3334  operand5_(operand5),
3335  operand6_(operand6),
3336  operand7_(operand7),
3337  operand8_(operand8)
3338  {}
3339 
3340  inline SeqWalkType init(IntVec const & extents,
3341  GhostData const & ghosts,
3342  IntVec const & hasBC) const {
3343  return SeqWalkType(functorFactory_,
3344  operand1_.init(extents, ghosts, hasBC),
3345  operand2_.init(extents, ghosts, hasBC),
3346  operand3_.init(extents, ghosts, hasBC),
3347  operand4_.init(extents, ghosts, hasBC),
3348  operand5_.init(extents, ghosts, hasBC),
3349  operand6_.init(extents, ghosts, hasBC),
3350  operand7_.init(extents, ghosts, hasBC),
3351  operand8_.init(extents, ghosts, hasBC));
3352  }
3353 
3354  private:
3355  FunctorFactory const & functorFactory_;
3356 
3357  Operand1 const operand1_;
3358 
3359  Operand2 const operand2_;
3360 
3361  Operand3 const operand3_;
3362 
3363  Operand4 const operand4_;
3364 
3365  Operand5 const operand5_;
3366 
3367  Operand6 const operand6_;
3368 
3369  Operand7 const operand7_;
3370 
3371  Operand8 const operand8_;
3372  }
3373  #endif
3374  /* ENABLE_THREADS */;
3375  template<typename FunctorFactory,
3376  typename Operand1,
3377  typename Operand2,
3378  typename Operand3,
3379  typename Operand4,
3380  typename Operand5,
3381  typename Operand6,
3382  typename Operand7,
3383  typename Operand8>
3384  struct ApplyPointwise8<SeqWalk,
3385  FunctorFactory,
3386  Operand1,
3387  Operand2,
3388  Operand3,
3389  Operand4,
3390  Operand5,
3391  Operand6,
3392  Operand7,
3393  Operand8> {
3394  public:
3395  typename Operand1::value_type typedef value_type;
3396 
3397  ApplyPointwise8(FunctorFactory const & functorFactory,
3398  Operand1 const & operand1,
3399  Operand2 const & operand2,
3400  Operand3 const & operand3,
3401  Operand4 const & operand4,
3402  Operand5 const & operand5,
3403  Operand6 const & operand6,
3404  Operand7 const & operand7,
3405  Operand8 const & operand8)
3406  : functor_(functorFactory()),
3407  operand1_(operand1),
3408  operand2_(operand2),
3409  operand3_(operand3),
3410  operand4_(operand4),
3411  operand5_(operand5),
3412  operand6_(operand6),
3413  operand7_(operand7),
3414  operand8_(operand8)
3415  {}
3416 
3417  inline value_type eval(int const x, int const y, int const z) const {
3418  return functor_(operand1_.eval(x, y, z),
3419  operand2_.eval(x, y, z),
3420  operand3_.eval(x, y, z),
3421  operand4_.eval(x, y, z),
3422  operand5_.eval(x, y, z),
3423  operand6_.eval(x, y, z),
3424  operand7_.eval(x, y, z),
3425  operand8_.eval(x, y, z));
3426  }
3427 
3428  private:
3429  typename FunctorFactory::Operator functor_;
3430 
3431  Operand1 operand1_;
3432 
3433  Operand2 operand2_;
3434 
3435  Operand3 operand3_;
3436 
3437  Operand4 operand4_;
3438 
3439  Operand5 operand5_;
3440 
3441  Operand6 operand6_;
3442 
3443  Operand7 operand7_;
3444 
3445  Operand8 operand8_;
3446  };
3447  #ifdef __CUDACC__
3448  template<typename FunctorFactory,
3449  typename Operand1,
3450  typename Operand2,
3451  typename Operand3,
3452  typename Operand4,
3453  typename Operand5,
3454  typename Operand6,
3455  typename Operand7,
3456  typename Operand8>
3457  struct ApplyPointwise8<GPUWalk,
3458  FunctorFactory,
3459  Operand1,
3460  Operand2,
3461  Operand3,
3462  Operand4,
3463  Operand5,
3464  Operand6,
3465  Operand7,
3466  Operand8> {
3467  public:
3468  typename Operand1::value_type typedef value_type;
3469 
3470  ApplyPointwise8(Operand1 const & operand1,
3471  Operand2 const & operand2,
3472  Operand3 const & operand3,
3473  Operand4 const & operand4,
3474  Operand5 const & operand5,
3475  Operand6 const & operand6,
3476  Operand7 const & operand7,
3477  Operand8 const & operand8)
3478  : operand1_(operand1),
3479  operand2_(operand2),
3480  operand3_(operand3),
3481  operand4_(operand4),
3482  operand5_(operand5),
3483  operand6_(operand6),
3484  operand7_(operand7),
3485  operand8_(operand8)
3486  {}
3487 
3488  __device__ inline value_type eval(int const x,
3489  int const y,
3490  int const z) const { return 0; }
3491 
3492  private:
3493  Operand1 operand1_;
3494 
3495  Operand2 operand2_;
3496 
3497  Operand3 operand3_;
3498 
3499  Operand4 operand4_;
3500 
3501  Operand5 operand5_;
3502 
3503  Operand6 operand6_;
3504 
3505  Operand7 operand7_;
3506 
3507  Operand8 operand8_;
3508  }
3509  #endif
3510  /* __CUDACC__ */;
3511 
3512  template<typename Operand1,
3513  typename Operand2,
3514  typename Operand3,
3515  typename Operand4,
3516  typename Operand5,
3517  typename Operand6,
3518  typename Operand7,
3519  typename Operand8>
3522  typename RefineFieldType<typename
3524  Result,
3525  typename
3526  RefineFieldType<typename
3528  Result,
3529  typename
3530  RefineFieldType<typename
3532  Result,
3533  typename
3534  RefineFieldType<typename
3536  Result,
3537  typename
3538  RefineFieldType<typename
3540  Result,
3541  typename
3542  RefineFieldType<typename
3544  Result,
3545  typename
3547  Result>::
3548  Result>::
3549  Result>::
3550  Result>::
3551  Result>::
3552  Result>::Result>::
3553  Result typedef Result;
3554  };
3555 
3556  template<typename FunctorFactory,
3557  typename Operand1,
3558  typename Operand2,
3559  typename Operand3,
3560  typename Operand4,
3561  typename Operand5,
3562  typename Operand6,
3563  typename Operand7,
3564  typename Operand8>
3565  inline NeboExpression<ApplyPointwise8<Initial,
3566  FunctorFactory,
3567  typename FinalType<Operand1>::
3568  Result,
3569  typename FinalType<Operand2>::
3570  Result,
3571  typename FinalType<Operand3>::
3572  Result,
3573  typename FinalType<Operand4>::
3574  Result,
3575  typename FinalType<Operand5>::
3576  Result,
3577  typename FinalType<Operand6>::
3578  Result,
3579  typename FinalType<Operand7>::
3580  Result,
3581  typename FinalType<Operand8>::
3582  Result>,
3583  typename FindFieldType8<Operand1,
3584  Operand2,
3585  Operand3,
3586  Operand4,
3587  Operand5,
3588  Operand6,
3589  Operand7,
3590  Operand8>::Result>
3591  apply_pointwise(FunctorFactory const & functorFactory,
3592  Operand1 operand1,
3593  Operand2 operand2,
3594  Operand3 operand3,
3595  Operand4 operand4,
3596  Operand5 operand5,
3597  Operand6 operand6,
3598  Operand7 operand7,
3599  Operand8 operand8) {
3600  ApplyPointwise8<Initial,
3601  FunctorFactory,
3602  typename FinalType<Operand1>::Result,
3603  typename FinalType<Operand2>::Result,
3604  typename FinalType<Operand3>::Result,
3605  typename FinalType<Operand4>::Result,
3606  typename FinalType<Operand5>::Result,
3607  typename FinalType<Operand6>::Result,
3608  typename FinalType<Operand7>::Result,
3609  typename FinalType<Operand8>::Result> typedef
3610  ReturnType;
3611 
3612  NeboExpression<ReturnType,
3613  typename FindFieldType8<Operand1,
3614  Operand2,
3615  Operand3,
3616  Operand4,
3617  Operand5,
3618  Operand6,
3619  Operand7,
3620  Operand8>::Result> typedef
3621  ReturnTerm;
3622 
3623  return ReturnTerm(ReturnType(functorFactory,
3624  normalize(operand1),
3625  normalize(operand2),
3626  normalize(operand3),
3627  normalize(operand4),
3628  normalize(operand5),
3629  normalize(operand6),
3630  normalize(operand7),
3631  normalize(operand8)));
3632  }
3633 
3634  template<typename CurrentMode,
3635  typename FunctorFactory,
3636  typename Operand1,
3637  typename Operand2,
3638  typename Operand3,
3639  typename Operand4,
3640  typename Operand5,
3641  typename Operand6,
3642  typename Operand7,
3643  typename Operand8,
3644  typename Operand9>
3646  template<typename FunctorFactory,
3647  typename Operand1,
3648  typename Operand2,
3649  typename Operand3,
3650  typename Operand4,
3651  typename Operand5,
3652  typename Operand6,
3653  typename Operand7,
3654  typename Operand8,
3655  typename Operand9>
3656  struct ApplyPointwise9<Initial,
3657  FunctorFactory,
3658  Operand1,
3659  Operand2,
3660  Operand3,
3661  Operand4,
3662  Operand5,
3663  Operand6,
3664  Operand7,
3665  Operand8,
3666  Operand9> {
3667  public:
3668  ApplyPointwise9<SeqWalk,
3669  FunctorFactory,
3670  typename Operand1::SeqWalkType,
3671  typename Operand2::SeqWalkType,
3672  typename Operand3::SeqWalkType,
3673  typename Operand4::SeqWalkType,
3674  typename Operand5::SeqWalkType,
3675  typename Operand6::SeqWalkType,
3676  typename Operand7::SeqWalkType,
3677  typename Operand8::SeqWalkType,
3678  typename Operand9::SeqWalkType> typedef SeqWalkType;
3679 
3680  #ifdef ENABLE_THREADS
3681  ApplyPointwise9<Resize,
3682  FunctorFactory,
3683  typename Operand1::ResizeType,
3684  typename Operand2::ResizeType,
3685  typename Operand3::ResizeType,
3686  typename Operand4::ResizeType,
3687  typename Operand5::ResizeType,
3688  typename Operand6::ResizeType,
3689  typename Operand7::ResizeType,
3690  typename Operand8::ResizeType,
3691  typename Operand9::ResizeType> typedef ResizeType;
3692  #endif
3693  /* ENABLE_THREADS */
3694 
3695  #ifdef __CUDACC__
3696  ApplyPointwise9<GPUWalk,
3697  FunctorFactory,
3698  typename Operand1::GPUWalkType,
3699  typename Operand2::GPUWalkType,
3700  typename Operand3::GPUWalkType,
3701  typename Operand4::GPUWalkType,
3702  typename Operand5::GPUWalkType,
3703  typename Operand6::GPUWalkType,
3704  typename Operand7::GPUWalkType,
3705  typename Operand8::GPUWalkType,
3706  typename Operand9::GPUWalkType> typedef GPUWalkType
3707  ;
3708  #endif
3709  /* __CUDACC__ */
3710 
3711  ApplyPointwise9(FunctorFactory const & functorFactory,
3712  Operand1 const & operand1,
3713  Operand2 const & operand2,
3714  Operand3 const & operand3,
3715  Operand4 const & operand4,
3716  Operand5 const & operand5,
3717  Operand6 const & operand6,
3718  Operand7 const & operand7,
3719  Operand8 const & operand8,
3720  Operand9 const & operand9)
3721  : functorFactory_(functorFactory),
3722  operand1_(operand1),
3723  operand2_(operand2),
3724  operand3_(operand3),
3725  operand4_(operand4),
3726  operand5_(operand5),
3727  operand6_(operand6),
3728  operand7_(operand7),
3729  operand8_(operand8),
3730  operand9_(operand9)
3731  {}
3732 
3733  inline GhostData ghosts_with_bc(void) const {
3734  return min(operand1_.ghosts_with_bc(),
3735  min(operand2_.ghosts_with_bc(),
3736  min(operand3_.ghosts_with_bc(),
3737  min(operand4_.ghosts_with_bc(),
3738  min(operand5_.ghosts_with_bc(),
3739  min(operand6_.ghosts_with_bc(),
3740  min(operand7_.ghosts_with_bc(),
3741  min(operand8_.ghosts_with_bc(),
3742  operand9_.ghosts_with_bc()))))))));
3743  }
3744 
3745  inline GhostData ghosts_without_bc(void) const {
3746  return min(operand1_.ghosts_without_bc(),
3747  min(operand2_.ghosts_without_bc(),
3748  min(operand3_.ghosts_without_bc(),
3749  min(operand4_.ghosts_without_bc(),
3750  min(operand5_.ghosts_without_bc(),
3751  min(operand6_.ghosts_without_bc(),
3752  min(operand7_.ghosts_without_bc(),
3753  min(operand8_.ghosts_without_bc(),
3754  operand9_.ghosts_without_bc()))))))));
3755  }
3756 
3757  inline bool has_extents(void) const {
3758  return (operand1_.has_extents() || operand2_.has_extents() ||
3759  operand3_.has_extents() || operand4_.has_extents() ||
3760  operand5_.has_extents() || operand6_.has_extents() ||
3761  operand7_.has_extents() || operand8_.has_extents() ||
3762  operand9_.has_extents());
3763  }
3764 
3765  inline IntVec extents(void) const {
3766  #ifndef NDEBUG
3767  if((operand1_.has_extents() || operand2_.has_extents() ||
3768  operand3_.has_extents() || operand4_.has_extents() ||
3769  operand5_.has_extents() || operand6_.has_extents() ||
3770  operand7_.has_extents() || operand8_.has_extents() ||
3771  operand9_.has_extents())) {
3772  IntVec extents;
3773 
3774  if(operand1_.has_extents()) { extents = operand1_.extents(); }
3775  else {
3776  if(operand2_.has_extents()) {
3777  extents = operand2_.extents();
3778  }
3779  else {
3780  if(operand3_.has_extents()) {
3781  extents = operand3_.extents();
3782  }
3783  else {
3784  if(operand4_.has_extents()) {
3785  extents = operand4_.extents();
3786  }
3787  else {
3788  if(operand5_.has_extents()) {
3789  extents = operand5_.extents();
3790  }
3791  else {
3792  if(operand6_.has_extents()) {
3793  extents = operand6_.extents();
3794  }
3795  else {
3796  if(operand7_.has_extents()) {
3797  extents = operand7_.extents();
3798  }
3799  else {
3800  if(operand8_.has_extents()) {
3801  extents = operand8_.extents();
3802  }
3803  else { extents = operand9_.extents(); };
3804  };
3805  };
3806  };
3807  };
3808  };
3809  };
3810  };
3811 
3812  if(operand1_.has_extents()) {
3813  assert(extents == operand1_.extents());
3814  };
3815 
3816  if(operand2_.has_extents()) {
3817  assert(extents == operand2_.extents());
3818  };
3819 
3820  if(operand3_.has_extents()) {
3821  assert(extents == operand3_.extents());
3822  };
3823 
3824  if(operand4_.has_extents()) {
3825  assert(extents == operand4_.extents());
3826  };
3827 
3828  if(operand5_.has_extents()) {
3829  assert(extents == operand5_.extents());
3830  };
3831 
3832  if(operand6_.has_extents()) {
3833  assert(extents == operand6_.extents());
3834  };
3835 
3836  if(operand7_.has_extents()) {
3837  assert(extents == operand7_.extents());
3838  };
3839 
3840  if(operand8_.has_extents()) {
3841  assert(extents == operand8_.extents());
3842  };
3843 
3844  if(operand9_.has_extents()) {
3845  assert(extents == operand9_.extents());
3846  };
3847  }
3848  #endif
3849  /* NDEBUG */;
3850 
3851  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
3852  ?
3853  operand2_.extents()
3854  : ((operand3_.has_extents()
3855  ?
3856  operand3_.extents()
3857  : ((operand4_.has_extents()
3858  ?
3859  operand4_.extents()
3860  :
3861  ((operand5_.has_extents()
3862  ?
3863  operand5_.extents()
3864  :
3865  ((operand6_.has_extents()
3866  ?
3867  operand6_.extents()
3868  :
3869  ((operand7_.has_extents()
3870  ?
3871  operand7_.extents()
3872  :
3873  ((operand8_.has_extents()
3874  ?
3875  operand8_.extents()
3876  :
3877  (operand9_.extents()))))))))))))))));
3878  }
3879 
3880  inline IntVec has_bc(void) const {
3881  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
3882  ?
3883  operand2_.has_bc()
3884  : ((operand3_.has_extents()
3885  ?
3886  operand3_.has_bc()
3887  : ((operand4_.has_extents()
3888  ?
3889  operand4_.has_bc()
3890  : ((operand5_.has_extents()
3891  ?
3892  operand5_.has_bc()
3893  :
3894  ((operand6_.has_extents()
3895  ?
3896  operand6_.has_bc()
3897  :
3898  ((operand7_.has_extents()
3899  ?
3900  operand7_.has_bc()
3901  :
3902  ((operand8_.has_extents()
3903  ?
3904  operand8_.has_bc()
3905  :
3906  (operand9_.has_bc()))))))))))))))));
3907  }
3908 
3909  inline SeqWalkType init(IntVec const & extents,
3910  GhostData const & ghosts,
3911  IntVec const & hasBC) const {
3912  return SeqWalkType(functorFactory_,
3913  operand1_.init(extents, ghosts, hasBC),
3914  operand2_.init(extents, ghosts, hasBC),
3915  operand3_.init(extents, ghosts, hasBC),
3916  operand4_.init(extents, ghosts, hasBC),
3917  operand5_.init(extents, ghosts, hasBC),
3918  operand6_.init(extents, ghosts, hasBC),
3919  operand7_.init(extents, ghosts, hasBC),
3920  operand8_.init(extents, ghosts, hasBC),
3921  operand9_.init(extents, ghosts, hasBC));
3922  }
3923 
3924  #ifdef ENABLE_THREADS
3925  inline ResizeType resize(void) const {
3926  return ResizeType(functorFactory_,
3927  operand1_.resize(),
3928  operand2_.resize(),
3929  operand3_.resize(),
3930  operand4_.resize(),
3931  operand5_.resize(),
3932  operand6_.resize(),
3933  operand7_.resize(),
3934  operand8_.resize(),
3935  operand9_.resize());
3936  }
3937  #endif
3938  /* ENABLE_THREADS */
3939 
3940  #ifdef __CUDACC__
3941  inline bool cpu_ready(void) const {
3942  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
3943  operand3_.cpu_ready() && operand4_.cpu_ready() &&
3944  operand5_.cpu_ready() && operand6_.cpu_ready() &&
3945  operand7_.cpu_ready() && operand8_.cpu_ready() &&
3946  operand9_.cpu_ready());
3947  }
3948 
3949  inline bool gpu_ready(int const deviceIndex) const { return false; }
3950 
3951  inline GPUWalkType gpu_init(IntVec const & extents,
3952  GhostData const & ghosts,
3953  IntVec const & hasBC,
3954  int const deviceIndex,
3955  cudaStream_t const & lhsStream) const {
3956  return GPUWalkType(operand1_.gpu_init(extents,
3957  ghosts,
3958  hasBC,
3959  deviceIndex,
3960  lhsStream),
3961  operand2_.gpu_init(extents,
3962  ghosts,
3963  hasBC,
3964  deviceIndex,
3965  lhsStream),
3966  operand3_.gpu_init(extents,
3967  ghosts,
3968  hasBC,
3969  deviceIndex,
3970  lhsStream),
3971  operand4_.gpu_init(extents,
3972  ghosts,
3973  hasBC,
3974  deviceIndex,
3975  lhsStream),
3976  operand5_.gpu_init(extents,
3977  ghosts,
3978  hasBC,
3979  deviceIndex,
3980  lhsStream),
3981  operand6_.gpu_init(extents,
3982  ghosts,
3983  hasBC,
3984  deviceIndex,
3985  lhsStream),
3986  operand7_.gpu_init(extents,
3987  ghosts,
3988  hasBC,
3989  deviceIndex,
3990  lhsStream),
3991  operand8_.gpu_init(extents,
3992  ghosts,
3993  hasBC,
3994  deviceIndex,
3995  lhsStream),
3996  operand9_.gpu_init(extents,
3997  ghosts,
3998  hasBC,
3999  deviceIndex,
4000  lhsStream));
4001  }
4002 
4003  inline void stream_wait_event(cudaEvent_t const & event) const {
4004  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
4005  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
4006  operand5_.stream_wait_event(event); operand6_.stream_wait_event(event);
4007  operand7_.stream_wait_event(event); operand8_.stream_wait_event(event);
4008  operand9_.stream_wait_event(event);
4009  }
4010 
4011  #ifdef NEBO_GPU_TEST
4012  inline void gpu_prep(int const deviceIndex) const {
4013  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
4014  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
4015  operand5_.gpu_prep(deviceIndex); operand6_.gpu_prep(deviceIndex);
4016  operand7_.gpu_prep(deviceIndex); operand8_.gpu_prep(deviceIndex);
4017  operand9_.gpu_prep(deviceIndex);
4018  }
4019  #endif
4020  /* NEBO_GPU_TEST */
4021  #endif
4022  /* __CUDACC__ */
4023 
4024  private:
4025  FunctorFactory const & functorFactory_;
4026 
4027  Operand1 const operand1_;
4028 
4029  Operand2 const operand2_;
4030 
4031  Operand3 const operand3_;
4032 
4033  Operand4 const operand4_;
4034 
4035  Operand5 const operand5_;
4036 
4037  Operand6 const operand6_;
4038 
4039  Operand7 const operand7_;
4040 
4041  Operand8 const operand8_;
4042 
4043  Operand9 const operand9_;
4044  };
4045  #ifdef ENABLE_THREADS
4046  template<typename FunctorFactory,
4047  typename Operand1,
4048  typename Operand2,
4049  typename Operand3,
4050  typename Operand4,
4051  typename Operand5,
4052  typename Operand6,
4053  typename Operand7,
4054  typename Operand8,
4055  typename Operand9>
4056  struct ApplyPointwise9<Resize,
4057  FunctorFactory,
4058  Operand1,
4059  Operand2,
4060  Operand3,
4061  Operand4,
4062  Operand5,
4063  Operand6,
4064  Operand7,
4065  Operand8,
4066  Operand9> {
4067  public:
4068  ApplyPointwise9<SeqWalk,
4069  FunctorFactory,
4070  typename Operand1::SeqWalkType,
4071  typename Operand2::SeqWalkType,
4072  typename Operand3::SeqWalkType,
4073  typename Operand4::SeqWalkType,
4074  typename Operand5::SeqWalkType,
4075  typename Operand6::SeqWalkType,
4076  typename Operand7::SeqWalkType,
4077  typename Operand8::SeqWalkType,
4078  typename Operand9::SeqWalkType> typedef SeqWalkType
4079  ;
4080 
4081  ApplyPointwise9(FunctorFactory const & functorFactory,
4082  Operand1 const & operand1,
4083  Operand2 const & operand2,
4084  Operand3 const & operand3,
4085  Operand4 const & operand4,
4086  Operand5 const & operand5,
4087  Operand6 const & operand6,
4088  Operand7 const & operand7,
4089  Operand8 const & operand8,
4090  Operand9 const & operand9)
4091  : functorFactory_(functorFactory),
4092  operand1_(operand1),
4093  operand2_(operand2),
4094  operand3_(operand3),
4095  operand4_(operand4),
4096  operand5_(operand5),
4097  operand6_(operand6),
4098  operand7_(operand7),
4099  operand8_(operand8),
4100  operand9_(operand9)
4101  {}
4102 
4103  inline SeqWalkType init(IntVec const & extents,
4104  GhostData const & ghosts,
4105  IntVec const & hasBC) const {
4106  return SeqWalkType(functorFactory_,
4107  operand1_.init(extents, ghosts, hasBC),
4108  operand2_.init(extents, ghosts, hasBC),
4109  operand3_.init(extents, ghosts, hasBC),
4110  operand4_.init(extents, ghosts, hasBC),
4111  operand5_.init(extents, ghosts, hasBC),
4112  operand6_.init(extents, ghosts, hasBC),
4113  operand7_.init(extents, ghosts, hasBC),
4114  operand8_.init(extents, ghosts, hasBC),
4115  operand9_.init(extents, ghosts, hasBC));
4116  }
4117 
4118  private:
4119  FunctorFactory const & functorFactory_;
4120 
4121  Operand1 const operand1_;
4122 
4123  Operand2 const operand2_;
4124 
4125  Operand3 const operand3_;
4126 
4127  Operand4 const operand4_;
4128 
4129  Operand5 const operand5_;
4130 
4131  Operand6 const operand6_;
4132 
4133  Operand7 const operand7_;
4134 
4135  Operand8 const operand8_;
4136 
4137  Operand9 const operand9_;
4138  }
4139  #endif
4140  /* ENABLE_THREADS */;
4141  template<typename FunctorFactory,
4142  typename Operand1,
4143  typename Operand2,
4144  typename Operand3,
4145  typename Operand4,
4146  typename Operand5,
4147  typename Operand6,
4148  typename Operand7,
4149  typename Operand8,
4150  typename Operand9>
4151  struct ApplyPointwise9<SeqWalk,
4152  FunctorFactory,
4153  Operand1,
4154  Operand2,
4155  Operand3,
4156  Operand4,
4157  Operand5,
4158  Operand6,
4159  Operand7,
4160  Operand8,
4161  Operand9> {
4162  public:
4163  typename Operand1::value_type typedef value_type;
4164 
4165  ApplyPointwise9(FunctorFactory const & functorFactory,
4166  Operand1 const & operand1,
4167  Operand2 const & operand2,
4168  Operand3 const & operand3,
4169  Operand4 const & operand4,
4170  Operand5 const & operand5,
4171  Operand6 const & operand6,
4172  Operand7 const & operand7,
4173  Operand8 const & operand8,
4174  Operand9 const & operand9)
4175  : functor_(functorFactory()),
4176  operand1_(operand1),
4177  operand2_(operand2),
4178  operand3_(operand3),
4179  operand4_(operand4),
4180  operand5_(operand5),
4181  operand6_(operand6),
4182  operand7_(operand7),
4183  operand8_(operand8),
4184  operand9_(operand9)
4185  {}
4186 
4187  inline value_type eval(int const x, int const y, int const z) const {
4188  return functor_(operand1_.eval(x, y, z),
4189  operand2_.eval(x, y, z),
4190  operand3_.eval(x, y, z),
4191  operand4_.eval(x, y, z),
4192  operand5_.eval(x, y, z),
4193  operand6_.eval(x, y, z),
4194  operand7_.eval(x, y, z),
4195  operand8_.eval(x, y, z),
4196  operand9_.eval(x, y, z));
4197  }
4198 
4199  private:
4200  typename FunctorFactory::Operator functor_;
4201 
4202  Operand1 operand1_;
4203 
4204  Operand2 operand2_;
4205 
4206  Operand3 operand3_;
4207 
4208  Operand4 operand4_;
4209 
4210  Operand5 operand5_;
4211 
4212  Operand6 operand6_;
4213 
4214  Operand7 operand7_;
4215 
4216  Operand8 operand8_;
4217 
4218  Operand9 operand9_;
4219  };
4220  #ifdef __CUDACC__
4221  template<typename FunctorFactory,
4222  typename Operand1,
4223  typename Operand2,
4224  typename Operand3,
4225  typename Operand4,
4226  typename Operand5,
4227  typename Operand6,
4228  typename Operand7,
4229  typename Operand8,
4230  typename Operand9>
4231  struct ApplyPointwise9<GPUWalk,
4232  FunctorFactory,
4233  Operand1,
4234  Operand2,
4235  Operand3,
4236  Operand4,
4237  Operand5,
4238  Operand6,
4239  Operand7,
4240  Operand8,
4241  Operand9> {
4242  public:
4243  typename Operand1::value_type typedef value_type;
4244 
4245  ApplyPointwise9(Operand1 const & operand1,
4246  Operand2 const & operand2,
4247  Operand3 const & operand3,
4248  Operand4 const & operand4,
4249  Operand5 const & operand5,
4250  Operand6 const & operand6,
4251  Operand7 const & operand7,
4252  Operand8 const & operand8,
4253  Operand9 const & operand9)
4254  : operand1_(operand1),
4255  operand2_(operand2),
4256  operand3_(operand3),
4257  operand4_(operand4),
4258  operand5_(operand5),
4259  operand6_(operand6),
4260  operand7_(operand7),
4261  operand8_(operand8),
4262  operand9_(operand9)
4263  {}
4264 
4265  __device__ inline value_type eval(int const x,
4266  int const y,
4267  int const z) const { return 0; }
4268 
4269  private:
4270  Operand1 operand1_;
4271 
4272  Operand2 operand2_;
4273 
4274  Operand3 operand3_;
4275 
4276  Operand4 operand4_;
4277 
4278  Operand5 operand5_;
4279 
4280  Operand6 operand6_;
4281 
4282  Operand7 operand7_;
4283 
4284  Operand8 operand8_;
4285 
4286  Operand9 operand9_;
4287  }
4288  #endif
4289  /* __CUDACC__ */;
4290 
4291  template<typename Operand1,
4292  typename Operand2,
4293  typename Operand3,
4294  typename Operand4,
4295  typename Operand5,
4296  typename Operand6,
4297  typename Operand7,
4298  typename Operand8,
4299  typename Operand9>
4302  typename RefineFieldType<typename
4304  Result,
4305  typename
4306  RefineFieldType<typename
4308  Result,
4309  typename
4310  RefineFieldType<typename
4312  Result,
4313  typename
4314  RefineFieldType<typename
4316  Result,
4317  typename
4318  RefineFieldType<typename
4320  Result,
4321  typename
4322  RefineFieldType<typename
4324  Result,
4325  typename
4326  RefineFieldType<typename
4328  Result,
4329  typename
4331  Result>::
4332  Result>::
4333  Result>::
4334  Result>::
4335  Result>::
4336  Result>::
4337  Result>::Result>::
4338  Result typedef Result;
4339  };
4340 
4341  template<typename FunctorFactory,
4342  typename Operand1,
4343  typename Operand2,
4344  typename Operand3,
4345  typename Operand4,
4346  typename Operand5,
4347  typename Operand6,
4348  typename Operand7,
4349  typename Operand8,
4350  typename Operand9>
4351  inline NeboExpression<ApplyPointwise9<Initial,
4352  FunctorFactory,
4353  typename FinalType<Operand1>::
4354  Result,
4355  typename FinalType<Operand2>::
4356  Result,
4357  typename FinalType<Operand3>::
4358  Result,
4359  typename FinalType<Operand4>::
4360  Result,
4361  typename FinalType<Operand5>::
4362  Result,
4363  typename FinalType<Operand6>::
4364  Result,
4365  typename FinalType<Operand7>::
4366  Result,
4367  typename FinalType<Operand8>::
4368  Result,
4369  typename FinalType<Operand9>::
4370  Result>,
4371  typename FindFieldType9<Operand1,
4372  Operand2,
4373  Operand3,
4374  Operand4,
4375  Operand5,
4376  Operand6,
4377  Operand7,
4378  Operand8,
4379  Operand9>::Result>
4380  apply_pointwise(FunctorFactory const & functorFactory,
4381  Operand1 operand1,
4382  Operand2 operand2,
4383  Operand3 operand3,
4384  Operand4 operand4,
4385  Operand5 operand5,
4386  Operand6 operand6,
4387  Operand7 operand7,
4388  Operand8 operand8,
4389  Operand9 operand9) {
4390  ApplyPointwise9<Initial,
4391  FunctorFactory,
4392  typename FinalType<Operand1>::Result,
4393  typename FinalType<Operand2>::Result,
4394  typename FinalType<Operand3>::Result,
4395  typename FinalType<Operand4>::Result,
4396  typename FinalType<Operand5>::Result,
4397  typename FinalType<Operand6>::Result,
4398  typename FinalType<Operand7>::Result,
4399  typename FinalType<Operand8>::Result,
4400  typename FinalType<Operand9>::Result> typedef
4401  ReturnType;
4402 
4403  NeboExpression<ReturnType,
4404  typename FindFieldType9<Operand1,
4405  Operand2,
4406  Operand3,
4407  Operand4,
4408  Operand5,
4409  Operand6,
4410  Operand7,
4411  Operand8,
4412  Operand9>::Result> typedef
4413  ReturnTerm;
4414 
4415  return ReturnTerm(ReturnType(functorFactory,
4416  normalize(operand1),
4417  normalize(operand2),
4418  normalize(operand3),
4419  normalize(operand4),
4420  normalize(operand5),
4421  normalize(operand6),
4422  normalize(operand7),
4423  normalize(operand8),
4424  normalize(operand9)));
4425  }
4426 
4427  template<typename CurrentMode,
4428  typename FunctorFactory,
4429  typename Operand1,
4430  typename Operand2,
4431  typename Operand3,
4432  typename Operand4,
4433  typename Operand5,
4434  typename Operand6,
4435  typename Operand7,
4436  typename Operand8,
4437  typename Operand9,
4438  typename Operand10>
4440  template<typename FunctorFactory,
4441  typename Operand1,
4442  typename Operand2,
4443  typename Operand3,
4444  typename Operand4,
4445  typename Operand5,
4446  typename Operand6,
4447  typename Operand7,
4448  typename Operand8,
4449  typename Operand9,
4450  typename Operand10>
4451  struct ApplyPointwise10<Initial,
4452  FunctorFactory,
4453  Operand1,
4454  Operand2,
4455  Operand3,
4456  Operand4,
4457  Operand5,
4458  Operand6,
4459  Operand7,
4460  Operand8,
4461  Operand9,
4462  Operand10> {
4463  public:
4464  ApplyPointwise10<SeqWalk,
4465  FunctorFactory,
4466  typename Operand1::SeqWalkType,
4467  typename Operand2::SeqWalkType,
4468  typename Operand3::SeqWalkType,
4469  typename Operand4::SeqWalkType,
4470  typename Operand5::SeqWalkType,
4471  typename Operand6::SeqWalkType,
4472  typename Operand7::SeqWalkType,
4473  typename Operand8::SeqWalkType,
4474  typename Operand9::SeqWalkType,
4475  typename Operand10::SeqWalkType> typedef SeqWalkType;
4476 
4477  #ifdef ENABLE_THREADS
4478  ApplyPointwise10<Resize,
4479  FunctorFactory,
4480  typename Operand1::ResizeType,
4481  typename Operand2::ResizeType,
4482  typename Operand3::ResizeType,
4483  typename Operand4::ResizeType,
4484  typename Operand5::ResizeType,
4485  typename Operand6::ResizeType,
4486  typename Operand7::ResizeType,
4487  typename Operand8::ResizeType,
4488  typename Operand9::ResizeType,
4489  typename Operand10::ResizeType> typedef ResizeType
4490  ;
4491  #endif
4492  /* ENABLE_THREADS */
4493 
4494  #ifdef __CUDACC__
4495  ApplyPointwise10<GPUWalk,
4496  FunctorFactory,
4497  typename Operand1::GPUWalkType,
4498  typename Operand2::GPUWalkType,
4499  typename Operand3::GPUWalkType,
4500  typename Operand4::GPUWalkType,
4501  typename Operand5::GPUWalkType,
4502  typename Operand6::GPUWalkType,
4503  typename Operand7::GPUWalkType,
4504  typename Operand8::GPUWalkType,
4505  typename Operand9::GPUWalkType,
4506  typename Operand10::GPUWalkType> typedef
4507  GPUWalkType;
4508  #endif
4509  /* __CUDACC__ */
4510 
4511  ApplyPointwise10(FunctorFactory const & functorFactory,
4512  Operand1 const & operand1,
4513  Operand2 const & operand2,
4514  Operand3 const & operand3,
4515  Operand4 const & operand4,
4516  Operand5 const & operand5,
4517  Operand6 const & operand6,
4518  Operand7 const & operand7,
4519  Operand8 const & operand8,
4520  Operand9 const & operand9,
4521  Operand10 const & operand10)
4522  : functorFactory_(functorFactory),
4523  operand1_(operand1),
4524  operand2_(operand2),
4525  operand3_(operand3),
4526  operand4_(operand4),
4527  operand5_(operand5),
4528  operand6_(operand6),
4529  operand7_(operand7),
4530  operand8_(operand8),
4531  operand9_(operand9),
4532  operand10_(operand10)
4533  {}
4534 
4535  inline GhostData ghosts_with_bc(void) const {
4536  return min(operand1_.ghosts_with_bc(),
4537  min(operand2_.ghosts_with_bc(),
4538  min(operand3_.ghosts_with_bc(),
4539  min(operand4_.ghosts_with_bc(),
4540  min(operand5_.ghosts_with_bc(),
4541  min(operand6_.ghosts_with_bc(),
4542  min(operand7_.ghosts_with_bc(),
4543  min(operand8_.ghosts_with_bc(),
4544  min(operand9_.ghosts_with_bc(),
4545  operand10_.ghosts_with_bc())))))))));
4546  }
4547 
4548  inline GhostData ghosts_without_bc(void) const {
4549  return min(operand1_.ghosts_without_bc(),
4550  min(operand2_.ghosts_without_bc(),
4551  min(operand3_.ghosts_without_bc(),
4552  min(operand4_.ghosts_without_bc(),
4553  min(operand5_.ghosts_without_bc(),
4554  min(operand6_.ghosts_without_bc(),
4555  min(operand7_.ghosts_without_bc(),
4556  min(operand8_.ghosts_without_bc(),
4557  min(operand9_.ghosts_without_bc(),
4558  operand10_.ghosts_without_bc())))))))));
4559  }
4560 
4561  inline bool has_extents(void) const {
4562  return (operand1_.has_extents() || operand2_.has_extents() ||
4563  operand3_.has_extents() || operand4_.has_extents() ||
4564  operand5_.has_extents() || operand6_.has_extents() ||
4565  operand7_.has_extents() || operand8_.has_extents() ||
4566  operand9_.has_extents() || operand10_.has_extents());
4567  }
4568 
4569  inline IntVec extents(void) const {
4570  #ifndef NDEBUG
4571  if((operand1_.has_extents() || operand2_.has_extents() ||
4572  operand3_.has_extents() || operand4_.has_extents() ||
4573  operand5_.has_extents() || operand6_.has_extents() ||
4574  operand7_.has_extents() || operand8_.has_extents() ||
4575  operand9_.has_extents() || operand10_.has_extents())) {
4576  IntVec extents;
4577 
4578  if(operand1_.has_extents()) { extents = operand1_.extents(); }
4579  else {
4580  if(operand2_.has_extents()) {
4581  extents = operand2_.extents();
4582  }
4583  else {
4584  if(operand3_.has_extents()) {
4585  extents = operand3_.extents();
4586  }
4587  else {
4588  if(operand4_.has_extents()) {
4589  extents = operand4_.extents();
4590  }
4591  else {
4592  if(operand5_.has_extents()) {
4593  extents = operand5_.extents();
4594  }
4595  else {
4596  if(operand6_.has_extents()) {
4597  extents = operand6_.extents();
4598  }
4599  else {
4600  if(operand7_.has_extents()) {
4601  extents = operand7_.extents();
4602  }
4603  else {
4604  if(operand8_.has_extents()) {
4605  extents = operand8_.extents();
4606  }
4607  else {
4608  if(operand9_.has_extents()) {
4609  extents = operand9_.extents();
4610  }
4611  else {
4612  extents = operand10_.extents();
4613  };
4614  };
4615  };
4616  };
4617  };
4618  };
4619  };
4620  };
4621  };
4622 
4623  if(operand1_.has_extents()) {
4624  assert(extents == operand1_.extents());
4625  };
4626 
4627  if(operand2_.has_extents()) {
4628  assert(extents == operand2_.extents());
4629  };
4630 
4631  if(operand3_.has_extents()) {
4632  assert(extents == operand3_.extents());
4633  };
4634 
4635  if(operand4_.has_extents()) {
4636  assert(extents == operand4_.extents());
4637  };
4638 
4639  if(operand5_.has_extents()) {
4640  assert(extents == operand5_.extents());
4641  };
4642 
4643  if(operand6_.has_extents()) {
4644  assert(extents == operand6_.extents());
4645  };
4646 
4647  if(operand7_.has_extents()) {
4648  assert(extents == operand7_.extents());
4649  };
4650 
4651  if(operand8_.has_extents()) {
4652  assert(extents == operand8_.extents());
4653  };
4654 
4655  if(operand9_.has_extents()) {
4656  assert(extents == operand9_.extents());
4657  };
4658 
4659  if(operand10_.has_extents()) {
4660  assert(extents == operand10_.extents());
4661  };
4662  }
4663  #endif
4664  /* NDEBUG */;
4665 
4666  return (operand1_.has_extents() ? operand1_.extents() : ((operand2_.has_extents()
4667  ?
4668  operand2_.extents()
4669  : ((operand3_.has_extents()
4670  ?
4671  operand3_.extents()
4672  : ((operand4_.has_extents()
4673  ?
4674  operand4_.extents()
4675  :
4676  ((operand5_.has_extents()
4677  ?
4678  operand5_.extents()
4679  :
4680  ((operand6_.has_extents()
4681  ?
4682  operand6_.extents()
4683  :
4684  ((operand7_.has_extents()
4685  ?
4686  operand7_.extents()
4687  :
4688  ((operand8_.has_extents()
4689  ?
4690  operand8_.extents()
4691  :
4692  ((operand9_.has_extents()
4693  ?
4694  operand9_.extents()
4695  :
4696  (operand10_.extents()))))))))))))))))));
4697  }
4698 
4699  inline IntVec has_bc(void) const {
4700  return (operand1_.has_extents() ? operand1_.has_bc() : ((operand2_.has_extents()
4701  ?
4702  operand2_.has_bc()
4703  : ((operand3_.has_extents()
4704  ?
4705  operand3_.has_bc()
4706  : ((operand4_.has_extents()
4707  ?
4708  operand4_.has_bc()
4709  : ((operand5_.has_extents()
4710  ?
4711  operand5_.has_bc()
4712  :
4713  ((operand6_.has_extents()
4714  ?
4715  operand6_.has_bc()
4716  :
4717  ((operand7_.has_extents()
4718  ?
4719  operand7_.has_bc()
4720  :
4721  ((operand8_.has_extents()
4722  ?
4723  operand8_.has_bc()
4724  :
4725  ((operand9_.has_extents()
4726  ?
4727  operand9_.has_bc()
4728  :
4729  (operand10_.has_bc()))))))))))))))))));
4730  }
4731 
4732  inline SeqWalkType init(IntVec const & extents,
4733  GhostData const & ghosts,
4734  IntVec const & hasBC) const {
4735  return SeqWalkType(functorFactory_,
4736  operand1_.init(extents, ghosts, hasBC),
4737  operand2_.init(extents, ghosts, hasBC),
4738  operand3_.init(extents, ghosts, hasBC),
4739  operand4_.init(extents, ghosts, hasBC),
4740  operand5_.init(extents, ghosts, hasBC),
4741  operand6_.init(extents, ghosts, hasBC),
4742  operand7_.init(extents, ghosts, hasBC),
4743  operand8_.init(extents, ghosts, hasBC),
4744  operand9_.init(extents, ghosts, hasBC),
4745  operand10_.init(extents, ghosts, hasBC));
4746  }
4747 
4748  #ifdef ENABLE_THREADS
4749  inline ResizeType resize(void) const {
4750  return ResizeType(functorFactory_,
4751  operand1_.resize(),
4752  operand2_.resize(),
4753  operand3_.resize(),
4754  operand4_.resize(),
4755  operand5_.resize(),
4756  operand6_.resize(),
4757  operand7_.resize(),
4758  operand8_.resize(),
4759  operand9_.resize(),
4760  operand10_.resize());
4761  }
4762  #endif
4763  /* ENABLE_THREADS */
4764 
4765  #ifdef __CUDACC__
4766  inline bool cpu_ready(void) const {
4767  return (operand1_.cpu_ready() && operand2_.cpu_ready() &&
4768  operand3_.cpu_ready() && operand4_.cpu_ready() &&
4769  operand5_.cpu_ready() && operand6_.cpu_ready() &&
4770  operand7_.cpu_ready() && operand8_.cpu_ready() &&
4771  operand9_.cpu_ready() && operand10_.cpu_ready());
4772  }
4773 
4774  inline bool gpu_ready(int const deviceIndex) const { return false; }
4775 
4776  inline GPUWalkType gpu_init(IntVec const & extents,
4777  GhostData const & ghosts,
4778  IntVec const & hasBC,
4779  int const deviceIndex,
4780  cudaStream_t const & lhsStream) const {
4781  return GPUWalkType(operand1_.gpu_init(extents,
4782  ghosts,
4783  hasBC,
4784  deviceIndex,
4785  lhsStream),
4786  operand2_.gpu_init(extents,
4787  ghosts,
4788  hasBC,
4789  deviceIndex,
4790  lhsStream),
4791  operand3_.gpu_init(extents,
4792  ghosts,
4793  hasBC,
4794  deviceIndex,
4795  lhsStream),
4796  operand4_.gpu_init(extents,
4797  ghosts,
4798  hasBC,
4799  deviceIndex,
4800  lhsStream),
4801  operand5_.gpu_init(extents,
4802  ghosts,
4803  hasBC,
4804  deviceIndex,
4805  lhsStream),
4806  operand6_.gpu_init(extents,
4807  ghosts,
4808  hasBC,
4809  deviceIndex,
4810  lhsStream),
4811  operand7_.gpu_init(extents,
4812  ghosts,
4813  hasBC,
4814  deviceIndex,
4815  lhsStream),
4816  operand8_.gpu_init(extents,
4817  ghosts,
4818  hasBC,
4819  deviceIndex,
4820  lhsStream),
4821  operand9_.gpu_init(extents,
4822  ghosts,
4823  hasBC,
4824  deviceIndex,
4825  lhsStream),
4826  operand10_.gpu_init(extents,
4827  ghosts,
4828  hasBC,
4829  deviceIndex,
4830  lhsStream));
4831  }
4832 
4833  inline void stream_wait_event(cudaEvent_t const & event) const {
4834  operand1_.stream_wait_event(event); operand2_.stream_wait_event(event);
4835  operand3_.stream_wait_event(event); operand4_.stream_wait_event(event);
4836  operand5_.stream_wait_event(event); operand6_.stream_wait_event(event);
4837  operand7_.stream_wait_event(event); operand8_.stream_wait_event(event);
4838  operand9_.stream_wait_event(event); operand10_.stream_wait_event(event);
4839  }
4840 
4841  #ifdef NEBO_GPU_TEST
4842  inline void gpu_prep(int const deviceIndex) const {
4843  operand1_.gpu_prep(deviceIndex); operand2_.gpu_prep(deviceIndex);
4844  operand3_.gpu_prep(deviceIndex); operand4_.gpu_prep(deviceIndex);
4845  operand5_.gpu_prep(deviceIndex); operand6_.gpu_prep(deviceIndex);
4846  operand7_.gpu_prep(deviceIndex); operand8_.gpu_prep(deviceIndex);
4847  operand9_.gpu_prep(deviceIndex); operand10_.gpu_prep(deviceIndex);
4848  }
4849  #endif
4850  /* NEBO_GPU_TEST */
4851  #endif
4852  /* __CUDACC__ */
4853 
4854  private:
4855  FunctorFactory const & functorFactory_;
4856 
4857  Operand1 const operand1_;
4858 
4859  Operand2 const operand2_;
4860 
4861  Operand3 const operand3_;
4862 
4863  Operand4 const operand4_;
4864 
4865  Operand5 const operand5_;
4866 
4867  Operand6 const operand6_;
4868 
4869  Operand7 const operand7_;
4870 
4871  Operand8 const operand8_;
4872 
4873  Operand9 const operand9_;
4874 
4875  Operand10 const operand10_;
4876  };
4877  #ifdef ENABLE_THREADS
4878  template<typename FunctorFactory,
4879  typename Operand1,
4880  typename Operand2,
4881  typename Operand3,
4882  typename Operand4,
4883  typename Operand5,
4884  typename Operand6,
4885  typename Operand7,
4886  typename Operand8,
4887  typename Operand9,
4888  typename Operand10>
4889  struct ApplyPointwise10<Resize,
4890  FunctorFactory,
4891  Operand1,
4892  Operand2,
4893  Operand3,
4894  Operand4,
4895  Operand5,
4896  Operand6,
4897  Operand7,
4898  Operand8,
4899  Operand9,
4900  Operand10> {
4901  public:
4902  ApplyPointwise10<SeqWalk,
4903  FunctorFactory,
4904  typename Operand1::SeqWalkType,
4905  typename Operand2::SeqWalkType,
4906  typename Operand3::SeqWalkType,
4907  typename Operand4::SeqWalkType,
4908  typename Operand5::SeqWalkType,
4909  typename Operand6::SeqWalkType,
4910  typename Operand7::SeqWalkType,
4911  typename Operand8::SeqWalkType,
4912  typename Operand9::SeqWalkType,
4913  typename Operand10::SeqWalkType> typedef
4914  SeqWalkType;
4915 
4916  ApplyPointwise10(FunctorFactory const & functorFactory,
4917  Operand1 const & operand1,
4918  Operand2 const & operand2,
4919  Operand3 const & operand3,
4920  Operand4 const & operand4,
4921  Operand5 const & operand5,
4922  Operand6 const & operand6,
4923  Operand7 const & operand7,
4924  Operand8 const & operand8,
4925  Operand9 const & operand9,
4926  Operand10 const & operand10)
4927  : functorFactory_(functorFactory),
4928  operand1_(operand1),
4929  operand2_(operand2),
4930  operand3_(operand3),
4931  operand4_(operand4),
4932  operand5_(operand5),
4933  operand6_(operand6),
4934  operand7_(operand7),
4935  operand8_(operand8),
4936  operand9_(operand9),
4937  operand10_(operand10)
4938  {}
4939 
4940  inline SeqWalkType init(IntVec const & extents,
4941  GhostData const & ghosts,
4942  IntVec const & hasBC) const {
4943  return SeqWalkType(functorFactory_,
4944  operand1_.init(extents, ghosts, hasBC),
4945  operand2_.init(extents, ghosts, hasBC),
4946  operand3_.init(extents, ghosts, hasBC),
4947  operand4_.init(extents, ghosts, hasBC),
4948  operand5_.init(extents, ghosts, hasBC),
4949  operand6_.init(extents, ghosts, hasBC),
4950  operand7_.init(extents, ghosts, hasBC),
4951  operand8_.init(extents, ghosts, hasBC),
4952  operand9_.init(extents, ghosts, hasBC),
4953  operand10_.init(extents, ghosts, hasBC));
4954  }
4955 
4956  private:
4957  FunctorFactory const & functorFactory_;
4958 
4959  Operand1 const operand1_;
4960 
4961  Operand2 const operand2_;
4962 
4963  Operand3 const operand3_;
4964 
4965  Operand4 const operand4_;
4966 
4967  Operand5 const operand5_;
4968 
4969  Operand6 const operand6_;
4970 
4971  Operand7 const operand7_;
4972 
4973  Operand8 const operand8_;
4974 
4975  Operand9 const operand9_;
4976 
4977  Operand10 const operand10_;
4978  }
4979  #endif
4980  /* ENABLE_THREADS */;
4981  template<typename FunctorFactory,
4982  typename Operand1,
4983  typename Operand2,
4984  typename Operand3,
4985  typename Operand4,
4986  typename Operand5,
4987  typename Operand6,
4988  typename Operand7,
4989  typename Operand8,
4990  typename Operand9,
4991  typename Operand10>
4992  struct ApplyPointwise10<SeqWalk,
4993  FunctorFactory,
4994  Operand1,
4995  Operand2,
4996  Operand3,
4997  Operand4,
4998  Operand5,
4999  Operand6,
5000  Operand7,
5001  Operand8,
5002  Operand9,
5003  Operand10> {
5004  public:
5005  typename Operand1::value_type typedef value_type;
5006 
5007  ApplyPointwise10(FunctorFactory const & functorFactory,
5008  Operand1 const & operand1,
5009  Operand2 const & operand2,
5010  Operand3 const & operand3,
5011  Operand4 const & operand4,
5012  Operand5 const & operand5,
5013  Operand6 const & operand6,
5014  Operand7 const & operand7,
5015  Operand8 const & operand8,
5016  Operand9 const & operand9,
5017  Operand10 const & operand10)
5018  : functor_(functorFactory()),
5019  operand1_(operand1),
5020  operand2_(operand2),
5021  operand3_(operand3),
5022  operand4_(operand4),
5023  operand5_(operand5),
5024  operand6_(operand6),
5025  operand7_(operand7),
5026  operand8_(operand8),
5027  operand9_(operand9),
5028  operand10_(operand10)
5029  {}
5030 
5031  inline value_type eval(int const x, int const y, int const z) const {
5032  return functor_(operand1_.eval(x, y, z),
5033  operand2_.eval(x, y, z),
5034  operand3_.eval(x, y, z),
5035  operand4_.eval(x, y, z),
5036  operand5_.eval(x, y, z),
5037  operand6_.eval(x, y, z),
5038  operand7_.eval(x, y, z),
5039  operand8_.eval(x, y, z),
5040  operand9_.eval(x, y, z),
5041  operand10_.eval(x, y, z));
5042  }
5043 
5044  private:
5045  typename FunctorFactory::Operator functor_;
5046 
5047  Operand1 operand1_;
5048 
5049  Operand2 operand2_;
5050 
5051  Operand3 operand3_;
5052 
5053  Operand4 operand4_;
5054 
5055  Operand5 operand5_;
5056 
5057  Operand6 operand6_;
5058 
5059  Operand7 operand7_;
5060 
5061  Operand8 operand8_;
5062 
5063  Operand9 operand9_;
5064 
5065  Operand10 operand10_;
5066  };
5067  #ifdef __CUDACC__
5068  template<typename FunctorFactory,
5069  typename Operand1,
5070  typename Operand2,
5071  typename Operand3,
5072  typename Operand4,
5073  typename Operand5,
5074  typename Operand6,
5075  typename Operand7,
5076  typename Operand8,
5077  typename Operand9,
5078  typename Operand10>
5079  struct ApplyPointwise10<GPUWalk,
5080  FunctorFactory,
5081  Operand1,
5082  Operand2,
5083  Operand3,
5084  Operand4,
5085  Operand5,
5086  Operand6,
5087  Operand7,
5088  Operand8,
5089  Operand9,
5090  Operand10> {
5091  public:
5092  typename Operand1::value_type typedef value_type;
5093 
5094  ApplyPointwise10(Operand1 const & operand1,
5095  Operand2 const & operand2,
5096  Operand3 const & operand3,
5097  Operand4 const & operand4,
5098  Operand5 const & operand5,
5099  Operand6 const & operand6,
5100  Operand7 const & operand7,
5101  Operand8 const & operand8,
5102  Operand9 const & operand9,
5103  Operand10 const & operand10)
5104  : operand1_(operand1),
5105  operand2_(operand2),
5106  operand3_(operand3),
5107  operand4_(operand4),
5108  operand5_(operand5),
5109  operand6_(operand6),
5110  operand7_(operand7),
5111  operand8_(operand8),
5112  operand9_(operand9),
5113  operand10_(operand10)
5114  {}
5115 
5116  __device__ inline value_type eval(int const x,
5117  int const y,
5118  int const z) const { return 0; }
5119 
5120  private:
5121  Operand1 operand1_;
5122 
5123  Operand2 operand2_;
5124 
5125  Operand3 operand3_;
5126 
5127  Operand4 operand4_;
5128 
5129  Operand5 operand5_;
5130 
5131  Operand6 operand6_;
5132 
5133  Operand7 operand7_;
5134 
5135  Operand8 operand8_;
5136 
5137  Operand9 operand9_;
5138 
5139  Operand10 operand10_;
5140  }
5141  #endif
5142  /* __CUDACC__ */;
5143 
5144  template<typename Operand1,
5145  typename Operand2,
5146  typename Operand3,
5147  typename Operand4,
5148  typename Operand5,
5149  typename Operand6,
5150  typename Operand7,
5151  typename Operand8,
5152  typename Operand9,
5153  typename Operand10>
5156  typename RefineFieldType<typename
5158  Result,
5159  typename
5160  RefineFieldType<typename
5162  Result,
5163  typename
5164  RefineFieldType<typename
5166  Result,
5167  typename
5168  RefineFieldType<typename
5170  Result,
5171  typename
5172  RefineFieldType<typename
5174  Result,
5175  typename
5176  RefineFieldType<typename
5178  Result,
5179  typename
5180  RefineFieldType<typename
5182  Result,
5183  typename
5184  RefineFieldType<typename
5186  Result,
5187  typename
5189  Result>::
5190  Result>::
5191  Result>::
5192  Result>::
5193  Result>::
5194  Result>::
5195  Result>::
5196  Result>::Result>::
5197  Result typedef Result;
5198  };
5199 
5200  template<typename FunctorFactory,
5201  typename Operand1,
5202  typename Operand2,
5203  typename Operand3,
5204  typename Operand4,
5205  typename Operand5,
5206  typename Operand6,
5207  typename Operand7,
5208  typename Operand8,
5209  typename Operand9,
5210  typename Operand10>
5211  inline NeboExpression<ApplyPointwise10<Initial,
5212  FunctorFactory,
5213  typename FinalType<Operand1>::
5214  Result,
5215  typename FinalType<Operand2>::
5216  Result,
5217  typename FinalType<Operand3>::
5218  Result,
5219  typename FinalType<Operand4>::
5220  Result,
5221  typename FinalType<Operand5>::
5222  Result,
5223  typename FinalType<Operand6>::
5224  Result,
5225  typename FinalType<Operand7>::
5226  Result,
5227  typename FinalType<Operand8>::
5228  Result,
5229  typename FinalType<Operand9>::
5230  Result,
5231  typename FinalType<Operand10>::
5232  Result>,
5233  typename FindFieldType10<Operand1,
5234  Operand2,
5235  Operand3,
5236  Operand4,
5237  Operand5,
5238  Operand6,
5239  Operand7,
5240  Operand8,
5241  Operand9,
5242  Operand10>::Result>
5243  apply_pointwise(FunctorFactory const & functorFactory,
5244  Operand1 operand1,
5245  Operand2 operand2,
5246  Operand3 operand3,
5247  Operand4 operand4,
5248  Operand5 operand5,
5249  Operand6 operand6,
5250  Operand7 operand7,
5251  Operand8 operand8,
5252  Operand9 operand9,
5253  Operand10 operand10) {
5254  ApplyPointwise10<Initial,
5255  FunctorFactory,
5256  typename FinalType<Operand1>::Result,
5257  typename FinalType<Operand2>::Result,
5258  typename FinalType<Operand3>::Result,
5259  typename FinalType<Operand4>::Result,
5260  typename FinalType<Operand5>::Result,
5261  typename FinalType<Operand6>::Result,
5262  typename FinalType<Operand7>::Result,
5263  typename FinalType<Operand8>::Result,
5264  typename FinalType<Operand9>::Result,
5265  typename FinalType<Operand10>::Result> typedef
5266  ReturnType;
5267 
5268  NeboExpression<ReturnType,
5269  typename FindFieldType10<Operand1,
5270  Operand2,
5271  Operand3,
5272  Operand4,
5273  Operand5,
5274  Operand6,
5275  Operand7,
5276  Operand8,
5277  Operand9,
5278  Operand10>::Result> typedef
5279  ReturnTerm;
5280 
5281  return ReturnTerm(ReturnType(functorFactory,
5282  normalize(operand1),
5283  normalize(operand2),
5284  normalize(operand3),
5285  normalize(operand4),
5286  normalize(operand5),
5287  normalize(operand6),
5288  normalize(operand7),
5289  normalize(operand8),
5290  normalize(operand9),
5291  normalize(operand10)));
5292  }
5293 
5294  template<typename OperatorT>
5296 
5297 
5298  public:
5299  OperatorT typedef Operator;
5300 
5301  FunctorFactory0() {}
5302 
5303  inline OperatorT operator()(void) const { return OperatorT(); }
5304  };
5305 
5306  template<typename Operator>
5307  inline FunctorFactory0<Operator> factory(void) {
5308  return FunctorFactory0<Operator>();
5309  }
5310 
5311  template<typename OperatorT, typename Arg1>
5313  private:
5314  Arg1 const & arg1_;
5315 
5316  public:
5317  OperatorT typedef Operator;
5318 
5319  FunctorFactory1(Arg1 const & arg1)
5320  : arg1_(arg1)
5321  {}
5322 
5323  inline OperatorT operator()(void) const { return OperatorT(arg1_); }
5324  };
5325 
5326  template<typename Operator, typename Arg1>
5327  inline FunctorFactory1<Operator, Arg1> factory(Arg1 const & arg1) {
5328  return FunctorFactory1<Operator, Arg1>(arg1);
5329  }
5330 
5331  template<typename OperatorT, typename Arg1, typename Arg2>
5333  private:
5334  Arg1 const & arg1_;
5335 
5336  Arg2 const & arg2_;
5337 
5338  public:
5339  OperatorT typedef Operator;
5340 
5341  FunctorFactory2(Arg1 const & arg1, Arg2 const & arg2)
5342  : arg1_(arg1), arg2_(arg2)
5343  {}
5344 
5345  inline OperatorT operator()(void) const {
5346  return OperatorT(arg1_, arg2_);
5347  }
5348  };
5349 
5350  template<typename Operator, typename Arg1, typename Arg2>
5351  inline FunctorFactory2<Operator, Arg1, Arg2> factory(Arg1 const & arg1,
5352  Arg2 const & arg2) {
5353  return FunctorFactory2<Operator, Arg1, Arg2>(arg1, arg2);
5354  }
5355 
5356  template<typename OperatorT, typename Arg1, typename Arg2, typename Arg3>
5358  private:
5359  Arg1 const & arg1_;
5360 
5361  Arg2 const & arg2_;
5362 
5363  Arg3 const & arg3_;
5364 
5365  public:
5366  OperatorT typedef Operator;
5367 
5368  FunctorFactory3(Arg1 const & arg1,
5369  Arg2 const & arg2,
5370  Arg3 const & arg3)
5371  : arg1_(arg1), arg2_(arg2), arg3_(arg3)
5372  {}
5373 
5374  inline OperatorT operator()(void) const {
5375  return OperatorT(arg1_, arg2_, arg3_);
5376  }
5377  };
5378 
5379  template<typename Operator, typename Arg1, typename Arg2, typename Arg3>
5380  inline FunctorFactory3<Operator, Arg1, Arg2, Arg3> factory(Arg1 const &
5381  arg1,
5382  Arg2 const &
5383  arg2,
5384  Arg3 const &
5385  arg3) {
5386  return FunctorFactory3<Operator, Arg1, Arg2, Arg3>(arg1, arg2, arg3);
5387  }
5388 
5389  template<typename OperatorT,
5390  typename Arg1,
5391  typename Arg2,
5392  typename Arg3,
5393  typename Arg4>
5395  private:
5396  Arg1 const & arg1_;
5397 
5398  Arg2 const & arg2_;
5399 
5400  Arg3 const & arg3_;
5401 
5402  Arg4 const & arg4_;
5403 
5404  public:
5405  OperatorT typedef Operator;
5406 
5407  FunctorFactory4(Arg1 const & arg1,
5408  Arg2 const & arg2,
5409  Arg3 const & arg3,
5410  Arg4 const & arg4)
5411  : arg1_(arg1), arg2_(arg2), arg3_(arg3), arg4_(arg4)
5412  {}
5413 
5414  inline OperatorT operator()(void) const {
5415  return OperatorT(arg1_, arg2_, arg3_, arg4_);
5416  }
5417  };
5418 
5419  template<typename Operator,
5420  typename Arg1,
5421  typename Arg2,
5422  typename Arg3,
5423  typename Arg4>
5425  const &
5426  arg1,
5427  Arg2
5428  const &
5429  arg2,
5430  Arg3
5431  const &
5432  arg3,
5433  Arg4
5434  const &
5435  arg4) {
5437  arg2,
5438  arg3,
5439  arg4);
5440  }
5441 
5442  template<typename OperatorT,
5443  typename Arg1,
5444  typename Arg2,
5445  typename Arg3,
5446  typename Arg4,
5447  typename Arg5>
5449  private:
5450  Arg1 const & arg1_;
5451 
5452  Arg2 const & arg2_;
5453 
5454  Arg3 const & arg3_;
5455 
5456  Arg4 const & arg4_;
5457 
5458  Arg5 const & arg5_;
5459 
5460  public:
5461  OperatorT typedef Operator;
5462 
5463  FunctorFactory5(Arg1 const & arg1,
5464  Arg2 const & arg2,
5465  Arg3 const & arg3,
5466  Arg4 const & arg4,
5467  Arg5 const & arg5)
5468  : arg1_(arg1), arg2_(arg2), arg3_(arg3), arg4_(arg4), arg5_(arg5)
5469  {}
5470 
5471  inline OperatorT operator()(void) const {
5472  return OperatorT(arg1_, arg2_, arg3_, arg4_, arg5_);
5473  }
5474  };
5475 
5476  template<typename Operator,
5477  typename Arg1,
5478  typename Arg2,
5479  typename Arg3,
5480  typename Arg4,
5481  typename Arg5>
5483  const
5484  &
5485  arg1,
5486  Arg2
5487  const
5488  &
5489  arg2,
5490  Arg3
5491  const
5492  &
5493  arg3,
5494  Arg4
5495  const
5496  &
5497  arg4,
5498  Arg5
5499  const
5500  &
5501  arg5) {
5503  arg2,
5504  arg3,
5505  arg4,
5506  arg5);
5507  }
5508 
5509  template<typename OperatorT,
5510  typename Arg1,
5511  typename Arg2,
5512  typename Arg3,
5513  typename Arg4,
5514  typename Arg5,
5515  typename Arg6>
5517  private:
5518  Arg1 const & arg1_;
5519 
5520  Arg2 const & arg2_;
5521 
5522  Arg3 const & arg3_;
5523 
5524  Arg4 const & arg4_;
5525 
5526  Arg5 const & arg5_;
5527 
5528  Arg6 const & arg6_;
5529 
5530  public:
5531  OperatorT typedef Operator;
5532 
5533  FunctorFactory6(Arg1 const & arg1,
5534  Arg2 const & arg2,
5535  Arg3 const & arg3,
5536  Arg4 const & arg4,
5537  Arg5 const & arg5,
5538  Arg6 const & arg6)
5539  : arg1_(arg1),
5540  arg2_(arg2),
5541  arg3_(arg3),
5542  arg4_(arg4),
5543  arg5_(arg5),
5544  arg6_(arg6)
5545  {}
5546 
5547  inline OperatorT operator()(void) const {
5548  return OperatorT(arg1_, arg2_, arg3_, arg4_, arg5_, arg6_);
5549  }
5550  };
5551 
5552  template<typename Operator,
5553  typename Arg1,
5554  typename Arg2,
5555  typename Arg3,
5556  typename Arg4,
5557  typename Arg5,
5558  typename Arg6>
5560  factory(Arg1 const & arg1,
5561  Arg2 const & arg2,
5562  Arg3 const & arg3,
5563  Arg4 const & arg4,
5564  Arg5 const & arg5,
5565  Arg6 const & arg6) {
5567  arg2,
5568  arg3,
5569  arg4,
5570  arg5,
5571  arg6);
5572  }
5573 
5574  template<typename OperatorT,
5575  typename Arg1,
5576  typename Arg2,
5577  typename Arg3,
5578  typename Arg4,
5579  typename Arg5,
5580  typename Arg6,
5581  typename Arg7>
5583  private:
5584  Arg1 const & arg1_;
5585 
5586  Arg2 const & arg2_;
5587 
5588  Arg3 const & arg3_;
5589 
5590  Arg4 const & arg4_;
5591 
5592  Arg5 const & arg5_;
5593 
5594  Arg6 const & arg6_;
5595 
5596  Arg7 const & arg7_;
5597 
5598  public:
5599  OperatorT typedef Operator;
5600 
5601  FunctorFactory7(Arg1 const & arg1,
5602  Arg2 const & arg2,
5603  Arg3 const & arg3,
5604  Arg4 const & arg4,
5605  Arg5 const & arg5,
5606  Arg6 const & arg6,
5607  Arg7 const & arg7)
5608  : arg1_(arg1),
5609  arg2_(arg2),
5610  arg3_(arg3),
5611  arg4_(arg4),
5612  arg5_(arg5),
5613  arg6_(arg6),
5614  arg7_(arg7)
5615  {}
5616 
5617  inline OperatorT operator()(void) const {
5618  return OperatorT(arg1_, arg2_, arg3_, arg4_, arg5_, arg6_, arg7_);
5619  }
5620  };
5621 
5622  template<typename Operator,
5623  typename Arg1,
5624  typename Arg2,
5625  typename Arg3,
5626  typename Arg4,
5627  typename Arg5,
5628  typename Arg6,
5629  typename Arg7>
5631  factory(Arg1 const & arg1,
5632  Arg2 const & arg2,
5633  Arg3 const & arg3,
5634  Arg4 const & arg4,
5635  Arg5 const & arg5,
5636  Arg6 const & arg6,
5637  Arg7 const & arg7) {
5638  return FunctorFactory7<Operator,
5639  Arg1,
5640  Arg2,
5641  Arg3,
5642  Arg4,
5643  Arg5,
5644  Arg6,
5645  Arg7>(arg1, arg2, arg3, arg4, arg5, arg6, arg7);
5646  }
5647 
5648  template<typename OperatorT,
5649  typename Arg1,
5650  typename Arg2,
5651  typename Arg3,
5652  typename Arg4,
5653  typename Arg5,
5654  typename Arg6,
5655  typename Arg7,
5656  typename Arg8>
5658  private:
5659  Arg1 const & arg1_;
5660 
5661  Arg2 const & arg2_;
5662 
5663  Arg3 const & arg3_;
5664 
5665  Arg4 const & arg4_;
5666 
5667  Arg5 const & arg5_;
5668 
5669  Arg6 const & arg6_;
5670 
5671  Arg7 const & arg7_;
5672 
5673  Arg8 const & arg8_;
5674 
5675  public:
5676  OperatorT typedef Operator;
5677 
5678  FunctorFactory8(Arg1 const & arg1,
5679  Arg2 const & arg2,
5680  Arg3 const & arg3,
5681  Arg4 const & arg4,
5682  Arg5 const & arg5,
5683  Arg6 const & arg6,
5684  Arg7 const & arg7,
5685  Arg8 const & arg8)
5686  : arg1_(arg1),
5687  arg2_(arg2),
5688  arg3_(arg3),
5689  arg4_(arg4),
5690  arg5_(arg5),
5691  arg6_(arg6),
5692  arg7_(arg7),
5693  arg8_(arg8)
5694  {}
5695 
5696  inline OperatorT operator()(void) const {
5697  return OperatorT(arg1_,
5698  arg2_,
5699  arg3_,
5700  arg4_,
5701  arg5_,
5702  arg6_,
5703  arg7_,
5704  arg8_);
5705  }
5706  };
5707 
5708  template<typename Operator,
5709  typename Arg1,
5710  typename Arg2,
5711  typename Arg3,
5712  typename Arg4,
5713  typename Arg5,
5714  typename Arg6,
5715  typename Arg7,
5716  typename Arg8>
5717  inline FunctorFactory8<Operator,
5718  Arg1,
5719  Arg2,
5720  Arg3,
5721  Arg4,
5722  Arg5,
5723  Arg6,
5724  Arg7,
5725  Arg8> factory(Arg1 const & arg1,
5726  Arg2 const & arg2,
5727  Arg3 const & arg3,
5728  Arg4 const & arg4,
5729  Arg5 const & arg5,
5730  Arg6 const & arg6,
5731  Arg7 const & arg7,
5732  Arg8 const & arg8) {
5733  return FunctorFactory8<Operator,
5734  Arg1,
5735  Arg2,
5736  Arg3,
5737  Arg4,
5738  Arg5,
5739  Arg6,
5740  Arg7,
5741  Arg8>(arg1,
5742  arg2,
5743  arg3,
5744  arg4,
5745  arg5,
5746  arg6,
5747  arg7,
5748  arg8);
5749  }
5750 
5751  template<typename OperatorT,
5752  typename Arg1,
5753  typename Arg2,
5754  typename Arg3,
5755  typename Arg4,
5756  typename Arg5,
5757  typename Arg6,
5758  typename Arg7,
5759  typename Arg8,
5760  typename Arg9>
5762  private:
5763  Arg1 const & arg1_;
5764 
5765  Arg2 const & arg2_;
5766 
5767  Arg3 const & arg3_;
5768 
5769  Arg4 const & arg4_;
5770 
5771  Arg5 const & arg5_;
5772 
5773  Arg6 const & arg6_;
5774 
5775  Arg7 const & arg7_;
5776 
5777  Arg8 const & arg8_;
5778 
5779  Arg9 const & arg9_;
5780 
5781  public:
5782  OperatorT typedef Operator;
5783 
5784  FunctorFactory9(Arg1 const & arg1,
5785  Arg2 const & arg2,
5786  Arg3 const & arg3,
5787  Arg4 const & arg4,
5788  Arg5 const & arg5,
5789  Arg6 const & arg6,
5790  Arg7 const & arg7,
5791  Arg8 const & arg8,
5792  Arg9 const & arg9)
5793  : arg1_(arg1),
5794  arg2_(arg2),
5795  arg3_(arg3),
5796  arg4_(arg4),
5797  arg5_(arg5),
5798  arg6_(arg6),
5799  arg7_(arg7),
5800  arg8_(arg8),
5801  arg9_(arg9)
5802  {}
5803 
5804  inline OperatorT operator()(void) const {
5805  return OperatorT(arg1_,
5806  arg2_,
5807  arg3_,
5808  arg4_,
5809  arg5_,
5810  arg6_,
5811  arg7_,
5812  arg8_,
5813  arg9_);
5814  }
5815  };
5816 
5817  template<typename Operator,
5818  typename Arg1,
5819  typename Arg2,
5820  typename Arg3,
5821  typename Arg4,
5822  typename Arg5,
5823  typename Arg6,
5824  typename Arg7,
5825  typename Arg8,
5826  typename Arg9>
5827  inline FunctorFactory9<Operator,
5828  Arg1,
5829  Arg2,
5830  Arg3,
5831  Arg4,
5832  Arg5,
5833  Arg6,
5834  Arg7,
5835  Arg8,
5836  Arg9> factory(Arg1 const & arg1,
5837  Arg2 const & arg2,
5838  Arg3 const & arg3,
5839  Arg4 const & arg4,
5840  Arg5 const & arg5,
5841  Arg6 const & arg6,
5842  Arg7 const & arg7,
5843  Arg8 const & arg8,
5844  Arg9 const & arg9) {
5845  return FunctorFactory9<Operator,
5846  Arg1,
5847  Arg2,
5848  Arg3,
5849  Arg4,
5850  Arg5,
5851  Arg6,
5852  Arg7,
5853  Arg8,
5854  Arg9>(arg1,
5855  arg2,
5856  arg3,
5857  arg4,
5858  arg5,
5859  arg6,
5860  arg7,
5861  arg8,
5862  arg9);
5863  }
5864 
5865  template<typename OperatorT,
5866  typename Arg1,
5867  typename Arg2,
5868  typename Arg3,
5869  typename Arg4,
5870  typename Arg5,
5871  typename Arg6,
5872  typename Arg7,
5873  typename Arg8,
5874  typename Arg9,
5875  typename Arg10>
5877  private:
5878  Arg1 const & arg1_;
5879 
5880  Arg2 const & arg2_;
5881 
5882  Arg3 const & arg3_;
5883 
5884  Arg4 const & arg4_;
5885 
5886  Arg5 const & arg5_;
5887 
5888  Arg6 const & arg6_;
5889 
5890  Arg7 const & arg7_;
5891 
5892  Arg8 const & arg8_;
5893 
5894  Arg9 const & arg9_;
5895 
5896  Arg10 const & arg10_;
5897 
5898  public:
5899  OperatorT typedef Operator;
5900 
5901  FunctorFactory10(Arg1 const & arg1,
5902  Arg2 const & arg2,
5903  Arg3 const & arg3,
5904  Arg4 const & arg4,
5905  Arg5 const & arg5,
5906  Arg6 const & arg6,
5907  Arg7 const & arg7,
5908  Arg8 const & arg8,
5909  Arg9 const & arg9,
5910  Arg10 const & arg10)
5911  : arg1_(arg1),
5912  arg2_(arg2),
5913  arg3_(arg3),
5914  arg4_(arg4),
5915  arg5_(arg5),
5916  arg6_(arg6),
5917  arg7_(arg7),
5918  arg8_(arg8),
5919  arg9_(arg9),
5920  arg10_(arg10)
5921  {}
5922 
5923  inline OperatorT operator()(void) const {
5924  return OperatorT(arg1_,
5925  arg2_,
5926  arg3_,
5927  arg4_,
5928  arg5_,
5929  arg6_,
5930  arg7_,
5931  arg8_,
5932  arg9_,
5933  arg10_);
5934  }
5935  };
5936 
5937  template<typename Operator,
5938  typename Arg1,
5939  typename Arg2,
5940  typename Arg3,
5941  typename Arg4,
5942  typename Arg5,
5943  typename Arg6,
5944  typename Arg7,
5945  typename Arg8,
5946  typename Arg9,
5947  typename Arg10>
5948  inline FunctorFactory10<Operator,
5949  Arg1,
5950  Arg2,
5951  Arg3,
5952  Arg4,
5953  Arg5,
5954  Arg6,
5955  Arg7,
5956  Arg8,
5957  Arg9,
5958  Arg10> factory(Arg1 const & arg1,
5959  Arg2 const & arg2,
5960  Arg3 const & arg3,
5961  Arg4 const & arg4,
5962  Arg5 const & arg5,
5963  Arg6 const & arg6,
5964  Arg7 const & arg7,
5965  Arg8 const & arg8,
5966  Arg9 const & arg9,
5967  Arg10 const & arg10) {
5968  return FunctorFactory10<Operator,
5969  Arg1,
5970  Arg2,
5971  Arg3,
5972  Arg4,
5973  Arg5,
5974  Arg6,
5975  Arg7,
5976  Arg8,
5977  Arg9,
5978  Arg10>(arg1,
5979  arg2,
5980  arg3,
5981  arg4,
5982  arg5,
5983  arg6,
5984  arg7,
5985  arg8,
5986  arg9,
5987  arg10);
5988  };
5989  } /* SpatialOps */
5990 
5991 #endif
5992 /* FULMAR_APPLY_POINTWISE_H */
Holds information about the number of ghost cells on each side of the domain.
Definition: GhostData.h:54