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