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