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