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