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