SpatialOps
NeboStencilBuilder.h
1 /*
2  * Copyright (c) 2014 The University of Utah
3  *
4  * Permission is hereby granted, free of charge, to any person obtaining a copy
5  * of this software and associated documentation files (the "Software"), to
6  * deal in the Software without restriction, including without limitation the
7  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
8  * sell copies of the Software, and to permit persons to whom the Software is
9  * furnished to do so, subject to the following conditions:
10  *
11  * The above copyright notice and this permission notice shall be included in
12  * all copies or substantial portions of the Software.
13  *
14  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19  * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
20  * IN THE SOFTWARE.
21  */
22 
23 #ifndef NEBO_STENCIL_BUILDER_H
24 #define NEBO_STENCIL_BUILDER_H
25 
26 //To define new stencils use the following four macros.
27 //NEBO_FIRST_POINT and NEBO_FIRST_IJK define a new Nebo stencil point collection.
28 //NEBO_ADD_POINT and NEBO_ADD_IJK add a stencil point to an existing Nebo stencil point collection.
29 //
30 //NEBO_FIRST_POINT and NEBO_ADD_POINT are designed to work together.
31 //Likewise, NEBO_FIRST_IJK and NEBO_ADD_IJK are designed to work together.
32 //They can be mixed, but it is not recommended.
33 //
34 //NEBO_FIRST_POINT and NEBO_ADD_POINT are designed to work with stencil points based off of field types.
35 //For examples, look at Stencil2Collection, Stencil4Collection, and FDStencilCollection definitions in this file.
36 //
37 //NEBO_FIRST_IJK and NEBO_ADD_IJK are designed to work on constant stencil points, whose shape does not change depending on field type.
38 //For examples, look at NullStencilCollection and the seven BoxFilter*StencilCollection definitions in this file.
39 
40 
41 
42 //Define a new stencil from an IndexTriplet
43 #define NEBO_FIRST_POINT_WO_TN(POINT) NeboStencilPointCollection< POINT, NeboNil >
44 #define NEBO_FIRST_POINT(POINT) typename NEBO_FIRST_POINT_WO_TN(POINT)
45 
46 //Add a point (IndexTriplet) to an existing stencil
47 #define NEBO_ADD_POINT(POINT) template AddPoint< POINT >::Result
48 
49 //Define a new stencil from three constant integers
50 #define NEBO_FIRST_IJK(X, Y, Z) NeboStencilPointCollection< IndexTriplet<X,Y,Z>, NeboNil >
51 
52 //Add a point (three constant integers) to an existing stencil
53 #define NEBO_ADD_IJK(X, Y, Z) AddPoint< IndexTriplet<X,Y,Z> >::Result
54 
55 namespace SpatialOps {
56 
61  template< typename List, typename IT >
62  struct ListSubtract
63  {
66  result;
67  };
68 
69  template< typename Point, typename IT >
71  {
73  result;
74  };
75 
85  template<typename OperatorType, typename PntCltnT, typename SrcFieldT, typename DestFieldT>
87  public:
88  typedef OperatorType type;
89  typedef PntCltnT PointCollectionType;
90  typedef SrcFieldT SrcFieldType;
91  typedef DestFieldT DestFieldType;
94 
95  // typedefs for when argument is a Nebo expression
96  template<typename Arg>
97  struct WithArg {
100  };
101 
102  // typedefs for when argument is a field
106 
111  : coefCollection_(coefs)
112  {}
113 
114  ~NeboStencilBuilder() {}
115 
119  CoefCollection const & coefs(void) const { return coefCollection_; };
120 
126  void apply_to_field( const SrcFieldType & src, DestFieldType & dest ) const {
127  dest <<= operator()(src);
128  }
129 
134  inline FieldResult operator ()( const SrcFieldType & src ) const {
135  return FieldResult(FieldStencil(FieldArg(src), coefs()));
136  }
137 
142  template<typename Arg>
143  inline typename WithArg<Arg>::Result
144  operator ()( const NeboExpression<Arg, SrcFieldType> & src ) const {
145  typedef typename WithArg<Arg>::Stencil Stencil;
146  typedef typename WithArg<Arg>::Result Result;
147  return Result(Stencil(src.expr(), coefs()));
148  }
149 
150  private:
151  const CoefCollection coefCollection_;
152  };
153 
154  template<typename OperatorType, typename SrcFieldType, typename DestFieldType>
156  // source field offset
157  typedef typename SrcFieldType::Location::Offset SrcOffset;
158  // destination field offset
159  typedef typename DestFieldType::Location::Offset DestOffset;
160  // low (first) stencil point location (relative to the destination point)
162  // high (second) stencil point location (relative to the destination point)
164  // collection of all stencil points in this stencil
165  typedef NEBO_FIRST_POINT(LowStPt)::NEBO_ADD_POINT(HighStPt) StPtCollection;
166  };
167 
168  template<typename OperatorType, typename SrcFieldType, typename DestFieldType>
170  template<bool Boolean, typename True, typename False>
171  struct TemplateIf;
172 
173  template<typename True, typename False>
174  struct TemplateIf<true, True, False> { True typedef result; };
175 
176  template<typename True, typename False>
177  struct TemplateIf<false, True, False> { False typedef result; };
178 
179  // source field offset
180  typedef typename SrcFieldType::Location::Offset SrcOffset;
181  // destination field offset
182  typedef typename DestFieldType::Location::Offset DestOffset;
183  // unit vectors
187  // first direction (unit vector)
188  typedef typename TemplateIf<((int)(SrcOffset::X) != (int)(DestOffset::X)),
189  XUnit,
190  YUnit>::result FirstDir;
191  // second direction (unit vector)
192  typedef typename TemplateIf<((int)(SrcOffset::Z) != (int)(DestOffset::Z)),
193  ZUnit,
194  YUnit>::result SecondDir;
195  // source offset in the first direction
197  // source offset in the second direction
199  // destination offset in the first direction
201  // destination offset in the second direction
203  // low value in the first direction
204  typedef typename GreaterThan<SrcInFirstDir,
205  DestInFirstDir>::result::Negate LoValInFirstDir;
206  // high value in the first direction
207  typedef typename LessThan<SrcInFirstDir,
208  DestInFirstDir>::result HiValInFirstDir;
209  // low value in the second direction
210  typedef typename GreaterThan<SrcInSecondDir,
211  DestInSecondDir>::result::Negate LoValInSecondDir;
212  // high value in the second direction
213  typedef typename LessThan<SrcInSecondDir,
214  DestInSecondDir>::result HiValInSecondDir;
215  // stencil point locations (relative to the destination point)
220  // collection of all stencil points in this stencil
221  typedef NEBO_FIRST_POINT(StPt1)::NEBO_ADD_POINT(StPt2)
222  ::NEBO_ADD_POINT(StPt3)::NEBO_ADD_POINT(StPt4) StPtCollection;
223  };
224 
225  template<typename OperatorType, typename SrcFieldType, typename DestFieldType>
227  typedef typename OperatorType::DirT DirT;
228  typedef typename UnitTriplet<DirT>::type DirVec;
229  typedef typename DirVec::Negate LowStPt;
230  typedef DirVec HighStPt;
231  typedef NEBO_FIRST_POINT(LowStPt)::NEBO_ADD_POINT(HighStPt) StPtCollection;
232  };
233 
243  template<typename OperatorType, typename PntCltnT, typename SrcFieldT, typename DestFieldT>
245  public:
246  typedef OperatorType type;
247  typedef PntCltnT PointCollectionType;
248  typedef SrcFieldT SrcFieldType;
249  typedef DestFieldT DestFieldType;
252 
253  // typedefs for when argument is a Nebo expression
254  template<typename Arg>
255  struct WithArg {
258  };
259 
260  // typedefs for when argument is a field
264 
269  : coefCollection_(coefs)
270  {}
271 
273 
277  CoefCollection const & coefs(void) const { return coefCollection_; };
278 
284  void apply_to_field( const SrcFieldType & src, DestFieldType & dest ) const {
285  dest <<= operator()(src);
286  }
287 
292  inline FieldResult operator ()( const SrcFieldType & src ) const {
293  return FieldResult(FieldStencil(FieldArg(src), coefs()));
294  }
295 
300  template<typename Arg>
301  inline typename WithArg<Arg>::Result
302  operator ()( const NeboExpression<Arg, SrcFieldType> & src ) const {
303  typedef typename WithArg<Arg>::Stencil Stencil;
304  typedef typename WithArg<Arg>::Result Result;
305  return Result(Stencil(src.expr(), coefs()));
306  }
307 
308  private:
309  const CoefCollection coefCollection_;
310  };
311 
320  template<typename PntCltnT, typename SrcFieldT, typename DestFieldT>
322  public:
323  typedef PntCltnT PointCollectionType;
324  typedef SrcFieldT SrcFieldType;
325  typedef DestFieldT DestFieldType;
326 
327  // typedefs for when argument is a Nebo expression
328  template<typename Arg>
329  struct WithArg {
332  };
333 
334  // typedefs for when argument is a field
338 
343  {}
344 
346 
352  void apply_to_field( const SrcFieldType & src, DestFieldType & dest ) const {
353  dest <<= operator()(src);
354  }
355 
360  inline FieldResult operator ()( const SrcFieldType & src ) const {
361  return FieldResult(FieldStencil(FieldArg(src)));
362  }
363 
368  template<typename Arg>
369  inline typename WithArg<Arg>::Result
370  operator ()( const NeboExpression<Arg, SrcFieldType> & src ) const {
371  typedef typename WithArg<Arg>::Stencil Stencil;
372  typedef typename WithArg<Arg>::Result Result;
373  return Result(Stencil(src.expr()));
374  }
375  };
376 
378  typedef NEBO_FIRST_IJK(0, 0, 0) StPtCollection;
379  };
380 
389  template<typename PntCltnT, typename SrcFieldT, typename DestFieldT>
391  public:
392  typedef PntCltnT PointCollectionType;
393  typedef SrcFieldT SrcFieldType;
394  typedef DestFieldT DestFieldType;
395 
396  // typedefs for when argument is a Nebo expression
397  template<typename Arg>
398  struct WithArg {
403  };
404 
405  // typedefs for when argument is a field
411 
416  {}
417 
419 
425  void apply_to_field( const SrcFieldType & src, DestFieldType & dest ) const {
426  dest <<= operator()(src);
427  }
428 
433  inline FieldResult operator ()( const SrcFieldType & src ) const {
434  return FieldResult(FieldAverage(FieldStencil(FieldArg(src)),
435  FieldScalar(1.0 / double(PointCollectionType::length))));
436  }
437 
442  template<typename Arg>
443  inline typename WithArg<Arg>::Result
444  operator ()( const NeboExpression<Arg, SrcFieldType> & src ) const {
445  typedef typename WithArg<Arg>::Stencil Stencil;
446  typedef typename WithArg<Arg>::Scalar Scalar;
447  typedef typename WithArg<Arg>::Average Average;
448  typedef typename WithArg<Arg>::Result Result;
449  return Result(Average(Stencil(src.expr()),
450  Scalar(1.0 / double(PointCollectionType::length))));
451  }
452  };
453 
455  typedef NEBO_FIRST_IJK(-1,-1,-1)::NEBO_ADD_IJK( 0,-1,-1)::NEBO_ADD_IJK( 1,-1,-1)
456  ::NEBO_ADD_IJK(-1, 0,-1)::NEBO_ADD_IJK( 0, 0,-1)::NEBO_ADD_IJK( 1, 0,-1)
457  ::NEBO_ADD_IJK(-1, 1,-1)::NEBO_ADD_IJK( 0, 1,-1)::NEBO_ADD_IJK( 1, 1,-1)
458  ::NEBO_ADD_IJK(-1,-1, 0)::NEBO_ADD_IJK( 0,-1, 0)::NEBO_ADD_IJK( 1,-1, 0)
459  ::NEBO_ADD_IJK(-1, 0, 0)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 1, 0, 0)
460  ::NEBO_ADD_IJK(-1, 1, 0)::NEBO_ADD_IJK( 0, 1, 0)::NEBO_ADD_IJK( 1, 1, 0)
461  ::NEBO_ADD_IJK(-1,-1, 1)::NEBO_ADD_IJK( 0,-1, 1)::NEBO_ADD_IJK( 1,-1, 1)
462  ::NEBO_ADD_IJK(-1, 0, 1)::NEBO_ADD_IJK( 0, 0, 1)::NEBO_ADD_IJK( 1, 0, 1)
463  ::NEBO_ADD_IJK(-1, 1, 1)::NEBO_ADD_IJK( 0, 1, 1)::NEBO_ADD_IJK( 1, 1, 1)
464  StPtCollection;
465  };
466 
468  typedef NEBO_FIRST_IJK(-1,-1, 0)::NEBO_ADD_IJK( 0,-1, 0)::NEBO_ADD_IJK( 1,-1, 0)
469  ::NEBO_ADD_IJK(-1, 0, 0)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 1, 0, 0)
470  ::NEBO_ADD_IJK(-1, 1, 0)::NEBO_ADD_IJK( 0, 1, 0)::NEBO_ADD_IJK( 1, 1, 0)
471  StPtCollection;
472  };
473 
475  typedef NEBO_FIRST_IJK(-1, 0,-1)::NEBO_ADD_IJK( 0, 0,-1)::NEBO_ADD_IJK( 1, 0,-1)
476  ::NEBO_ADD_IJK(-1, 0, 0)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 1, 0, 0)
477  ::NEBO_ADD_IJK(-1, 0, 1)::NEBO_ADD_IJK( 0, 0, 1)::NEBO_ADD_IJK( 1, 0, 1)
478  StPtCollection;
479  };
480 
482  typedef NEBO_FIRST_IJK( 0,-1,-1)::NEBO_ADD_IJK( 0, 0,-1)::NEBO_ADD_IJK( 0, 1,-1)
483  ::NEBO_ADD_IJK( 0,-1, 0)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 0, 1, 0)
484  ::NEBO_ADD_IJK( 0,-1, 1)::NEBO_ADD_IJK( 0, 0, 1)::NEBO_ADD_IJK( 0, 1, 1)
485  StPtCollection;
486  };
487 
489  typedef NEBO_FIRST_IJK(-1, 0, 0)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 1, 0, 0)
490  StPtCollection;
491  };
492 
494  typedef NEBO_FIRST_IJK( 0,-1, 0)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 0, 1, 0)
495  StPtCollection;
496  };
497 
499  typedef NEBO_FIRST_IJK( 0, 0,-1)::NEBO_ADD_IJK( 0, 0, 0)::NEBO_ADD_IJK( 0, 0, 1)
500  StPtCollection;
501  };
502 
503  template<typename OperatorType, typename SrcFieldType, typename DestFieldType>
505  // source field offset
506  typedef typename SrcFieldType::Location::Offset SrcOffset;
507  // destination field offset
508  typedef typename DestFieldType::Location::Offset DestOffset;
509  // minus-side stencil point location
511  // plus-side stencil point location
512  typedef typename GreaterThan<SrcOffset,
513  DestOffset>::result::Negate PlusPoint;
514  };
515 
516  //must be a finite difference (FD) stencil, pull direction and points from operator:
517  template<typename OperatorType, typename FieldType>
518  struct MaskShiftPoints<OperatorType, FieldType, FieldType> {
519  typedef typename OperatorType::type OpT;
520  // FDStencilCollection:
522  // minus-side stencil point location
523  typedef typename FDStencil::DirVec MinusPoint;
524  // plus-side stencil point location
525  typedef typename FDStencil::DirVec::Negate PlusPoint;
526  };
527 
536  template<typename OperatorT, typename SrcFieldT, typename DestFieldT>
538  public:
539  typedef OperatorT OperatorType;
540  typedef SrcFieldT SrcFieldType;
541  typedef DestFieldT DestFieldType;
542 
545 
548 
550 
553 
556 
561  {}
562 
564 
569  inline MinusResult minus( const SrcMask & src ) const {
570  return MinusResult(MinusShift(Mask(src)));
571  }
572 
577  inline PlusResult plus( const SrcMask & src ) const {
578  return PlusResult(PlusShift(Mask(src)));
579  }
580  };
581 
585  template<typename CurrentMode,typename Arg0, typename Arg1,typename FieldType>
586  struct NeboSwitch;
587  template<typename Arg0, typename Arg1, typename FieldType>
588  struct NeboSwitch<Initial, Arg0, Arg1, FieldType> {
589  FieldType typedef field_type;
591 #ifdef ENABLE_THREADS
593  typedef ResizeType;
594 #endif
595  NeboSwitch(Arg0 const& arg0, Arg1 const& arg1, int which): arg0_(arg0), arg1_(arg1), which_(which){}
596 #define INVOKE_ARG_METHOD(name) (which_ == 0?arg0_.name:arg1_.name)
597  inline GhostData ghosts_with_bc(void) const {
598  if(which_ < 2)
599  return INVOKE_ARG_METHOD(ghosts_with_bc());
600  else
601  return max(arg0_.ghosts_with_bc(), arg1_.ghosts_with_bc());
602  }
603  inline GhostData ghosts_without_bc(void) const {
604  if(which_ < 2)
605  return INVOKE_ARG_METHOD(ghosts_with_bc());
606  else
607  return max(arg0_.ghosts_without_bc(), arg1_.ghosts_without_bc());
608  }
609  inline bool has_extents(void) const { return INVOKE_ARG_METHOD(has_extents()); }
610  inline IntVec extents(void) const { return INVOKE_ARG_METHOD(extents()); }
611  inline IntVec has_bc(void) const { return INVOKE_ARG_METHOD(has_bc());}
612  inline SeqWalkType init(IntVec const & extents, GhostData const & ghosts, IntVec const & hasBC) const {
613  return SeqWalkType(arg0_.init(extents, ghosts, hasBC),
614  arg1_.init(extents, ghosts, hasBC),
615  which_,
616  get_range(arg0_.extents(), arg0_.ghosts_without_bc()),
617  get_range(arg1_.extents(), arg1_.ghosts_without_bc()));
618  }
619 #ifdef ENABLE_THREADS
620  inline ResizeType resize(void) const {
621  return ResizeType(arg0_.resize(),
622  arg1_.resize(),
623  which_,
624  get_range(arg0_.extents(), arg0_.ghosts_without_bc()),
625  get_range(arg1_.extents(), arg1_.ghosts_without_bc()));
626  }
627 #endif
628 #ifdef __CUDACC__
630  inline bool cpu_ready(void) const { return INVOKE_ARG_METHOD(cpu_ready()); }
631  inline bool gpu_ready(int const deviceIndex) const {return INVOKE_ARG_METHOD(gpu_ready(deviceIndex));}
632  inline GPUWalkType gpu_init(IntVec const & extents,GhostData const & ghosts,IntVec const & hasBC, int const deviceIndex, cudaStream_t const & lhsStream) const {
633  return GPUWalkType(arg0_.gpu_init(extents, ghosts, hasBC, deviceIndex, lhsStream),
634  arg1_.gpu_init(extents, ghosts, hasBC, deviceIndex, lhsStream),
635  which_,
636  get_range(arg0_.extents(), arg0_.ghosts_without_bc()),
637  get_range(arg1_.extents(), arg1_.ghosts_without_bc()));
638  }
639  inline void stream_wait_event(cudaEvent_t const & event) const {
640  arg0_.stream_wait_event(event); arg1_.stream_wait_event(event);
641  }
642 # ifdef NEBO_GPU_TEST
643  inline void gpu_prep(int const deviceIndex) const {
644  arg0_.gpu_prep(deviceIndex);
645  arg1_.gpu_prep(deviceIndex);
646  }
647 # endif
648 #endif
649  private:
650  Arg0 const arg0_;
651  Arg1 const arg1_;
652  int which_;
653  static inline GhostData get_range(IntVec const& extents, GhostData const& ghosts)
654  {
655  return GhostData(-ghosts.get_minus(0),
656  extents[0] + ghosts.get_plus(0),
657  -ghosts.get_minus(1),
658  extents[1] + ghosts.get_plus(1),
659  -ghosts.get_minus(2),
660  extents[2] + ghosts.get_plus(2));
661  }
662  };
663 #ifdef ENABLE_THREADS
664  template<typename Arg0, typename Arg1, typename FieldType>
665  struct NeboSwitch<Resize, Arg0, Arg1, FieldType> {
666  FieldType typedef field_type;
668  NeboSwitch(Arg0 const & arg0, Arg1 const & arg1, int which, GhostData const& range0, GhostData const& range1)
669  : arg0_(arg0), arg1_(arg1), range0_(range0), range1_(range1), which_(which)
670  {}
671  inline SeqWalkType init(IntVec const & extents,GhostData const & ghosts, IntVec const & hasBC) const {
672  return SeqWalkType(arg0_.init(extents, ghosts, hasBC),
673  arg1_.init(extents, ghosts, hasBC),
674  which_,
675  _overlap(range0_, get_range(extents, ghosts)),
676  _overlap(range1_, get_range(extents, ghosts)));
677  }
678  private:
679  Arg0 const arg0_;
680  Arg1 const arg1_;
681  GhostData const range0_;
682  GhostData const range1_;
683  int which_;
684  static inline GhostData _overlap(const GhostData& g1, const GhostData& g2)
685  {
686  return GhostData(std::max(g1.get_minus(0), g2.get_minus(0)),
687  std::min(g1.get_plus(0), g2.get_plus(0)),
688  std::max(g1.get_minus(1), g2.get_minus(1)),
689  std::min(g1.get_plus(1), g2.get_plus(1)),
690  std::max(g1.get_minus(2), g2.get_minus(2)),
691  std::min(g1.get_plus(2), g2.get_plus(2)));
692  }
693  static inline GhostData get_range(IntVec const& extents, GhostData const& ghosts)
694  {
695  return GhostData(-ghosts.get_minus(0),
696  extents[0] + ghosts.get_plus(0),
697  -ghosts.get_minus(1),
698  extents[1] + ghosts.get_plus(1),
699  -ghosts.get_minus(2),
700  extents[2] + ghosts.get_plus(2));
701  }
702  };
703 #endif
704  template<typename Arg0, typename Arg1, typename FieldType>
705  struct NeboSwitch<SeqWalk, Arg0, Arg1, FieldType> {
706  FieldType typedef field_type;
707  NeboSwitch(Arg0 const & arg0, Arg1 const & arg1, int which, GhostData const& range0, GhostData const& range1)
708  : arg0_(arg0), arg1_(arg1), range0_(range0), range1_(range1), which_(which)
709  {}
710  inline bool eval(int const x, int const y, int const z) const {
711  if(which_ < 2)
712  {
713  return INVOKE_ARG_METHOD(eval(x,y,z));
714  }
715  else
716  {
717  return ((range0_.get_minus(0) <= x && x < range0_.get_plus(0) &&
718  range0_.get_minus(1) <= y && y < range0_.get_plus(1) &&
719  range0_.get_minus(2) <= z && z < range0_.get_plus(2))?arg0_.eval(x,y,z):0) ||
720  ((range1_.get_minus(0) <= x && x < range1_.get_plus(0) &&
721  range1_.get_minus(1) <= y && y < range1_.get_plus(1) &&
722  range1_.get_minus(2) <= z && z < range1_.get_plus(2))?arg1_.eval(x,y,z):0);
723  }
724  }
725  private:
726  Arg0 const arg0_;
727  Arg1 const arg1_;
728  GhostData const range0_;
729  GhostData const range1_;
730  int which_;
731  };
732 #ifdef __CUDACC__
733  template<typename Arg0, typename Arg1, typename FieldType>
734  struct NeboSwitch<GPUWalk, Arg0, Arg1, FieldType> {
735  FieldType typedef field_type;
736  NeboSwitch(Arg0 const & arg0, Arg1 const & arg1, int which, GhostData const& range0, GhostData const& range1)
737  : arg0_(arg0), arg1_(arg1), which_(which)
738  {
739  for(int i = 0; i < 3; i ++)
740  {
741  minus0[i] = range0.get_minus(i);
742  plus0[i] = range0.get_plus(i);
743  minus1[i] = range1.get_minus(i);
744  plus1[i] = range1.get_plus(i);
745  }
746  }
747  __device__ inline bool eval(int const x, int const y, int const z) const {
748  if(which_ < 2)
749  {
750  return INVOKE_ARG_METHOD(eval(x,y,z));
751  }
752  else
753  {
754  return ((minus0[0] <= x && x < plus0[0] &&
755  minus0[1] <= y && y < plus0[1] &&
756  minus0[2] <= z && z < plus0[2])?arg0_.eval(x,y,z):0) ||
757  ((minus1[0] <= x && x < plus1[0] &&
758  minus1[1] <= y && y < plus1[1] &&
759  minus1[2] <= z && z < plus1[2])?arg1_.eval(x,y,z):0);
760  }
761  }
762  private:
763  Arg0 const arg0_;
764  Arg1 const arg1_;
765  int minus0[3];
766  int plus0[3];
767  int minus1[3];
768  int plus1[3];
769  int which_;
770  };
771 #endif
772 
778  template <typename SrcMaskType, typename DestMaskType, typename Dir = NODIR>
780  public:
781  typedef std::vector<IntVec> Points;
782  typedef SrcMaskType FromType;
783  typedef DestMaskType ToType;
784 
786  struct type{
787  typedef Dir DirT;
788  };
789  };
790 
792  typename FromType::field_type,
793  typename ToType::field_type> Shift;
794 
795  typedef typename Shift::MinusPoint MinusPointT;
796  typedef typename Shift::PlusPoint PlusPointT;
797 
798  typedef typename Shift::MinusShift MinusShiftT;
799  typedef typename Shift::PlusShift PlusShiftT;
800 
802 
804 
805  ConvertedMask(const SrcMaskType& src, BCSide side1, BCSide side2, BCSide side3):src_(src), face_(side1)
806  {
807  if(side1 == MINUS_SIDE) neg_face_ = PLUS_SIDE;
808  else neg_face_ = MINUS_SIDE;
809  if(side2 == NO_SIDE) side2 = PLUS_SIDE;
810  if(side3 == NO_SIDE) side3 = PLUS_SIDE;
811  all_ = ((side2 != -1 && side2 != side1) || (side3 != -1 && side3 != side1));
812  }
813  inline BCSide which_face()
814  {
815  return neg_face_;
816  }
817  inline const Points& points() const
818  {
819  return src_.points();
820  }
821  inline const SrcMaskType& src_mask() const{
822  return src_;
823  }
824  inline Result mask() const{
825  return Result(Selector(MinusShiftT(src_), PlusShiftT(src_), all_?2:face_!= MINUS_SIDE));
826  }
827  /* this function make a spatial mask from a converted mask */
828  inline DestMaskType to_mask(const typename DestMaskType::field_type& field) const{
829  std::vector<IntVec> points;
830  for(std::vector<IntVec>::const_iterator it = src_.points().begin();
831  it != src_.points().end();
832  it ++)
833  {
834  if(all_)
835  {
836  _add_point(*it - MinusPointT::int_vec(), points, field.window_without_ghost(), field.get_ghost_data());
837  _add_point(*it - PlusPointT::int_vec(), points, field.window_without_ghost(), field.get_ghost_data());
838  }
839  else if(face_ == MINUS_SIDE)
840  {
841  _add_point(*it - MinusPointT::int_vec(), points, field.window_without_ghost(), field.get_ghost_data());
842  }
843  else
844  {
845  _add_point(*it - PlusPointT::int_vec(), points, field.window_without_ghost(), field.get_ghost_data());
846  }
847  }
848  return DestMaskType(field, points);
849  }
850  inline bool all() const{ return all_ == 1; }
851  inline operator Result() const { return mask(); }
852  inline IntVec get_shift_vec(int idx) const
853  {
854  if(all_ && idx == 0) return MinusPointT::int_vec();
855  else if(all_ && idx == 1) return PlusPointT::int_vec();
856  else if(face_ == MINUS_SIDE) return MinusPointT::int_vec();
857  else return PlusPointT::int_vec();
858  }
859  template <typename FieldType, typename Point>
861  {
864  return ReturnType(ShiftType(mask().expr()));
865  }
866  private:
867  //Shift shift_;
868  const SrcMaskType& src_;
869  BCSide neg_face_, face_;
870  bool all_;
871  static inline void _add_point(const IntVec& point, std::vector<IntVec>& vec, const MemoryWindow& maskWindow_, const GhostData& ghosts_)
872  {
873 
874  if(!(point[0] >= -ghosts_.get_minus(0)) ||
875  !(point[0] < (signed int)(maskWindow_.extent(0)) + ghosts_.get_plus(0)) ||
876  !(point[1] >= -ghosts_.get_minus(1)) ||
877  !(point[1] < (signed int)(maskWindow_.extent(1)) + ghosts_.get_plus(1)) ||
878  !(point[2] >= -ghosts_.get_minus(2)) ||
879  !(point[2] < (signed int)(maskWindow_.extent(2)) + ghosts_.get_plus(2))) return;
880  vec.push_back(point);
881  }
882  };
883  template<typename OpT, typename SrcFieldT, typename DestFieldT>
884  struct GetOperatorDirImp{typedef NODIR result;};
885  template<typename OpT, typename SrcFieldT>
886  struct GetOperatorDirImp<OpT, SrcFieldT, SrcFieldT>{typedef typename OpT::type::DirT result;};
887  template<typename OpT>
890  };
891 
892  template <typename Dir, typename SrcFieldT, typename DestFieldT>
893  struct GetFDOperatorDir{typedef NODIR result;};
894  template <typename Dir, typename SrcFieldT>
895  struct GetFDOperatorDir<Dir, SrcFieldT, SrcFieldT>{typedef Dir result;};
896 
905  template<typename OperatorT>
907  public:
908  typedef OperatorT OperatorType;
909  typedef typename OperatorType::SrcFieldType PhiFieldType;
910  typedef typename OperatorType::DestFieldType GammaFieldType;
911 
912  typedef typename OperatorType::PointCollectionType PointCollection;
913 
914  typedef typename PointCollection::Last LowPoint;
915  typedef typename PointCollection::AllButLast NonLowPoints;
916  typedef typename PointCollection::First HighPoint;
917  typedef typename PointCollection::AllButFirst NonHighPoints;
918 
919  typedef typename Subtract<IndexTriplet<0,0,0>, LowPoint>:: result LowGammaPoint;
920  typedef typename Subtract<IndexTriplet<0,0,0>, HighPoint>::result HighGammaPoint;
921  typedef typename ListSubtract<NonLowPoints, LowPoint>:: result NonLowSrcPoints;
924 
926 
931 
932  typedef std::vector<IntVec> Points;
933  typedef Points::const_iterator PointIterator;
934 
936 
940  NeboBoundaryConditionBuilder(OperatorType const & op)
941  : lowCoef_(op.coefs().last()),
942  highCoef_(op.coefs().coef()),
943  minusGamma_(1.0),
944  plusGamma_(1.0),
945  minusPhi_(op.coefs().all_but_last()),
946  plusPhi_(op.coefs().others()),
947  shift_()
948  {}
949 
951 
960  template<typename ExprType, typename MaskType>
961  inline void apply_imp(MaskType mask,
962  PhiFieldType& phi,
963  const IntVec& shift,
965  BCSide side) const
966  {
967  if(side == MINUS_SIDE)
968  cpu_apply<ExprType, MinusGammaType, MinusPhiType>(mask.points(),
969  shift,
970  minusGamma_,
971  minusPhi_,
972  phi,
973  gamma,
974  lowCoef_);
975  else
976  cpu_apply<ExprType, PlusGammaType, PlusPhiType>(mask.points(),
977  shift,
978  plusGamma_,
979  plusPhi_,
980  phi,
981  gamma,
982  highCoef_);
983 
984  }
991  template<typename ExprType>
992  inline void operator()(ConvertedPhiMask mask,
993  PhiFieldType& phi,
994  const NeboExpression<ExprType, GammaFieldType>& gamma) const
995  {
996  if(mask.all())
997  throw std::runtime_error("can not apply boundary condition on more than one face");
998  if(phi.active_device_index() == CPU_INDEX) {
999  apply_imp(mask,
1000  phi,
1001  IntVec(0,0,0),
1002  gamma,
1003  mask.which_face());
1004  } else {
1005  if(mask.which_face() == MINUS_SIDE)
1006  phi <<= cond(mask.src_mask(), (minusGamma_(gamma) - minusPhi_(phi)) / lowCoef_)
1007  (phi);
1008  else
1009  phi <<= cond(mask.src_mask(), (plusGamma_(gamma) - plusPhi_(phi)) / highCoef_)
1010  (phi);
1011  }
1012  }
1013  template<typename MaskType, typename GammaType>
1014  inline void apply_indirect_conversion(MaskType& mask, PhiFieldType& phi, GammaType gamma, bool minus, int idx) const
1015  {
1016  IntVec conv_shift = mask.get_shift_vec(idx);
1017  if(minus)
1018  apply_imp(mask,
1019  phi,
1020  Shift::MinusPoint::int_vec() + conv_shift,
1021  gamma,
1022  MINUS_SIDE);
1023  else
1024  apply_imp(mask,
1025  phi,
1026  Shift::PlusPoint::int_vec() + conv_shift,
1027  gamma,
1028  PLUS_SIDE);
1029  }
1030  template<typename ExprType, typename SrcMaskType, typename Dir>
1031  inline void operator()(ConvertedMask<const SpatialMask<SrcMaskType>, const SpatialMask<GammaFieldType>, Dir> mask,
1032  PhiFieldType& phi,
1034  bool minus) const
1035  {
1036  if(phi.active_device_index() == CPU_INDEX)
1037  {
1038  apply_indirect_conversion(mask, phi, gamma, minus, 0);
1039  if(mask.all()) apply_indirect_conversion(mask, phi, gamma, minus, 1);
1040  }
1041  else
1042  {
1043  if(minus)
1044  phi <<= cond(mask.template shift<PhiFieldType, typename Shift::MinusPoint>(), (minusGamma_(gamma) - minusPhi_(phi)) / lowCoef_)
1045  (phi);
1046  else
1047  phi <<= cond(mask.template shift<PhiFieldType, typename Shift::PlusPoint>(), (plusGamma_(gamma) - plusPhi_(phi)) / highCoef_)
1048  (phi);
1049  }
1050 
1051  }
1059  template<typename ExprType>
1060  inline void operator()(SpatialMask<GammaFieldType> mask,
1061  PhiFieldType & phi,
1063  bool minus) const {
1064  if(phi.active_device_index() == CPU_INDEX) {
1065  if(minus)
1066  apply_imp(mask,
1067  phi,
1068  Shift::MinusPoint::int_vec(),
1069  gamma,
1070  MINUS_SIDE);
1071  else
1072  apply_imp(mask,
1073  phi,
1074  Shift::PlusPoint::int_vec(),
1075  gamma,
1076  PLUS_SIDE);
1077  }
1078  else {
1079  if(minus)
1080  phi <<= cond(shift_.minus(mask), (minusGamma_(gamma) - minusPhi_(phi)) / lowCoef_)
1081  (phi);
1082  else
1083  phi <<= cond(shift_.plus(mask), (plusGamma_(gamma) - plusPhi_(phi)) / highCoef_)
1084  (phi);
1085  }
1086  }
1087 
1095  inline void operator()(SpatialMask<GammaFieldType> mask,
1096  PhiFieldType & phi,
1097  const GammaFieldType & gamma,
1098  bool minus) const
1099  {
1100  typedef NeboConstField<Initial, GammaFieldType> GammaField;
1102  (*this)(mask, phi, GammaExpr(GammaField(gamma)), minus);
1103  }
1104 
1111  inline void operator()(ConvertedPhiMask mask,
1112  PhiFieldType & phi,
1113  const GammaFieldType & gamma) const
1114  {
1115  typedef NeboConstField<Initial, GammaFieldType> GammaField;
1117  (*this)(mask, phi, GammaExpr(GammaField(gamma)));
1118  }
1119 
1120  template<typename SrcMaskType, typename Dir>
1121  inline void operator()(ConvertedMask<const SpatialMask<SrcMaskType>, const SpatialMask<GammaFieldType>, Dir> mask,
1122  PhiFieldType& phi,
1123  const GammaFieldType& gamma,
1124  bool minus) const
1125  {
1126  typedef NeboConstField<Initial, GammaFieldType> GammaField;
1128  (*this)(mask, phi, GammaExpr(GammaField(gamma)), minus);
1129  }
1137  inline void operator()(const SpatialMask<GammaFieldType> mask,
1138  PhiFieldType & phi,
1139  const double gamma,
1140  bool minus) const
1141  {
1142  typedef NeboScalar<Initial, double> GammaScalar;
1144  (*this)(mask, phi, GammaExpr(GammaScalar(gamma)), minus);
1145  }
1146 
1154  inline void operator()(ConvertedPhiMask mask,
1155  PhiFieldType & phi,
1156  const double gamma) const
1157  {
1158  typedef NeboScalar<Initial, double> GammaScalar;
1160  (*this)(mask, phi, GammaExpr(GammaScalar(gamma)));
1161  }
1162 
1163  template<typename SrcMaskType, typename Dir>
1164  inline void operator()(ConvertedMask<const SpatialMask<SrcMaskType>, const SpatialMask<GammaFieldType>, Dir> mask,
1165  PhiFieldType& phi,
1166  const double gamma,
1167  bool minus) const
1168  {
1169  typedef NeboScalar<Initial, double> GammaScalar;
1171  (*this)(mask, phi, GammaExpr(GammaScalar(gamma)), minus);
1172  }
1173 
1174  private:
1175 
1186  template<typename ExprType, typename ShiftGamma, typename ShiftPhi>
1187  inline void cpu_apply( const Points & points,
1188  const IntVec & shift,
1189  const ShiftGamma & shiftGamma,
1190  const ShiftPhi & shiftPhi,
1191  PhiFieldType & phi,
1193  const double coef ) const
1194  {
1195  typedef NeboField<Initial, PhiFieldType> LhsTypeInit;
1196  typedef typename LhsTypeInit::SeqWalkType LhsType;
1197  LhsTypeInit lhsInit(phi);
1198  LhsType lhs = lhsInit.init();
1199 
1200  typedef typename ShiftGamma::template WithArg<ExprType>::Stencil GammaType;
1201  typedef typename ShiftPhi::FieldStencil PhiType;
1202  typedef DiffOp<Initial, GammaType, PhiType> DiffType;
1203  typedef NeboScalar<Initial, double> NumType;
1204  typedef DivOp<Initial, DiffType, NumType> RhsTypeInit;
1205  typedef typename RhsTypeInit::SeqWalkType RhsType;
1206  //arguments to init() call are meaningless, but they are not used at all...
1207  RhsType rhs = RhsTypeInit(DiffType(shiftGamma(gamma).expr(),
1208  shiftPhi(phi).expr()),
1209  NumType(coef)).init(IntVec(0,0,0),
1210  GhostData(0),
1211  IntVec(0,0,0));
1212 
1213  PointIterator ip = points.begin();
1214  PointIterator const ep = points.end();
1215  for(; ip != ep; ip++) {
1216  const int x = (*ip)[0] - shift[0];
1217  const int y = (*ip)[1] - shift[1];
1218  const int z = (*ip)[2] - shift[2];
1219  lhs.ref(x, y, z) = rhs.eval(x, y, z);
1220  }
1221  }
1222 
1223  const double lowCoef_;
1224  const double highCoef_;
1225  const MinusGammaType minusGamma_;
1226  const PlusGammaType plusGamma_;
1227  const MinusPhiType minusPhi_;
1228  const PlusPhiType plusPhi_;
1229  const Shift shift_;
1230  };
1231  template <typename DestFieldLocation, typename SrcFieldLocation>
1233  template <int x, int y>
1234  struct Equals{
1235  enum {result = (x == y)?1:0};
1236  };
1237  typedef typename DestFieldLocation::Offset Dest_Offset;
1238  typedef typename SrcFieldLocation::Offset Src_Offset;
1242  };
1244  {
1245  if(Shift && face == NO_SIDE)
1246  throw std::runtime_error("trying to make invalid mask conversion");
1247  }
1248  };
1249  template <typename DstFieldType, typename SrcFieldType> struct MaskConversionFieldTypeCheck;
1250 #define ALLOW_MASK_CONVERSION(from, to) template <> struct MaskConversionFieldTypeCheck<to, from>{};
1251  ALLOW_MASK_CONVERSION(SVolField, XVolField)
1252  ALLOW_MASK_CONVERSION(SVolField, YVolField)
1253  ALLOW_MASK_CONVERSION(SVolField, ZVolField)
1254 
1255 #define ALLOW_MASK_CONVERSION_GROUP(dir)\
1256  ALLOW_MASK_CONVERSION(dir##VolField, dir##SurfXField)\
1257  ALLOW_MASK_CONVERSION(dir##VolField, dir##SurfYField)\
1258  ALLOW_MASK_CONVERSION(dir##VolField, dir##SurfZField)\
1259 
1260  ALLOW_MASK_CONVERSION_GROUP(X)
1261  ALLOW_MASK_CONVERSION_GROUP(Y)
1262  ALLOW_MASK_CONVERSION_GROUP(Z)
1263  ALLOW_MASK_CONVERSION_GROUP(S)
1264 
1272  template <typename DstFieldType, typename Dir, typename SrcMskType>
1273  static inline ConvertedMask<const SrcMskType,
1274  const SpatialMask<DstFieldType>,
1275  typename GetFDOperatorDir<Dir,
1276  DstFieldType,
1277  typename SrcMskType::field_type>::result
1278  > convert(const SrcMskType& src, const BCSide face1, const BCSide face2 = (BCSide)-1, const BCSide face3 = (BCSide)-1)
1279  {
1280  typedef typename SrcMskType::field_type SrcFieldType;
1281  typedef SpatialMask<DstFieldType> DstMskType;
1282  typedef typename DstFieldType::Location DstLocation;
1283  typedef typename SrcFieldType::Location SrcLocation;
1285 
1286  MaskConversionCheck<DstLocation, SrcLocation> conversion_check(face1);
1287  return ConvertedMask<const SrcMskType, const DstMskType, FDDir>(src, face1, face2, face3);
1288  }
1289 
1290  template <typename DstFieldType, typename SrcMskType>
1291  static inline ConvertedMask<const SrcMskType,
1293  typename GetFDOperatorDir<NODIR,
1294  DstFieldType,
1295  typename SrcMskType::field_type>::result
1296  > convert(const SrcMskType& src, const BCSide face1, const BCSide face2 = (BCSide)-1, const BCSide face3 = (BCSide)-1)
1297  {
1298  return convert<DstFieldType, NODIR, SrcMskType>(src, face1, face2, face3);
1299  }
1300 
1301  //TODO: move this to other files
1302  template<typename SrcT, typename DestT, typename Dir>
1303  static inline void masked_assign(
1304  const ConvertedMask<const SpatialMask<SrcT>,
1305  const SpatialMask<DestT>,
1306  Dir>& mask,
1307  DestT& field,
1308  const typename DestT::value_type& rhs)
1309  {
1310 #if __CUDACC__
1311  if(field.active_device_index() == CPU_INDEX)
1312 #endif
1313  masked_assign_imp(mask, field, NeboScalar<Initial, typename DestT::value_type>(rhs));
1314 #if __CUDACC__
1315  else
1316  field <<= cond(mask.mask(), rhs)(field);
1317 #endif
1318  }
1319 
1320  template<typename SrcT, typename DestT, typename Dir>
1321  static inline void masked_assign(
1322  const ConvertedMask<const SpatialMask<SrcT>,
1323  const SpatialMask<DestT>,
1324  Dir>& mask,
1325  DestT& field,
1326  const DestT& rhs)
1327  {
1328 #if __CUDACC__
1329  if(field.active_device_index() == CPU_INDEX)
1330 #endif
1331  masked_assign_imp(mask, field, NeboConstField<Initial, DestT>(rhs));
1332 #if __CUDACC__
1333  else
1334  field <<= cond(mask.mask(), rhs)(field);
1335 #endif
1336  }
1337 
1338  template<typename SrcT, typename DestT, typename Dir, typename RSHType, typename Expr>
1339  static inline void masked_assign(
1340  const ConvertedMask<const SpatialMask<SrcT>,
1341  const SpatialMask<DestT>,
1342  Dir>& mask,
1343  DestT& field,
1345  {
1346 #if __CUDACC__
1347  if(field.active_device_index() == CPU_INDEX)
1348 #endif
1349  masked_assign_imp(mask, field, rhs.expr());
1350 #if __CUDACC__
1351  else
1352  field <<= cond(mask.mask(), rhs)(field);
1353 #endif
1354 
1355  }
1356 
1357  template<typename SrcT, typename DestT, typename Dir, typename RHS>
1358  static inline void masked_assign_imp(
1359  const ConvertedMask<const SpatialMask<SrcT>,
1360  const SpatialMask<DestT>,
1361  Dir>& mask,
1362  DestT& field,
1363  const RHS& rhs)
1364  {
1365  typedef NeboField<Initial, DestT> LHSType;
1366  LHSType ilhs(field);
1367 #if __CUDACC__
1368  if(IS_GPU_INDEX(field.active_device_index())) {
1369  std::ostringstream msg;
1370  msg << "Nebo error in " << "Nebo Masked Assignment" << ":\n"
1371  ;
1372  msg << "Left-hand side of masked assignment allocated on ";
1373  msg << "GPU and this backend does not support GPU execution"
1374  ;
1375  msg << "\n";
1376  msg << "\t - " << __FILE__ << " : " << __LINE__;
1377  throw(std::runtime_error(msg.str()));
1378  }
1379  else {
1380  if(IS_CPU_INDEX(field.active_device_index())) {
1381  if(rhs.cpu_ready()) {
1382 #endif
1383  typename LHSType::SeqWalkType lhs = ilhs.init();
1384  typename RHS::SeqWalkType expr = rhs.init(IntVec(0,0,0),
1385  GhostData(0),
1386  IntVec(0,0,0));
1387 
1388 
1389  for(int i = 0; i <= mask.all(); i ++)
1390  {
1391  IntVec shift = mask.get_shift_vec(i);
1392 
1393  std::vector<IntVec>::const_iterator ip = mask.points().begin();
1394 
1395  std::vector<IntVec>::const_iterator const ep = mask.points().end();
1396 
1397  for(; ip != ep; ip++) {
1398  int const x = (*ip)[0] - shift[0];
1399 
1400  int const y = (*ip)[1] - shift[1];
1401 
1402  int const z = (*ip)[2] - shift[2];
1403 
1404  lhs.ref(x, y, z) = expr.eval(x, y, z);
1405  };
1406  }
1407 #if __CUDACC__
1408  }
1409  else {
1410  std::ostringstream msg;
1411  msg << "Nebo error in " << "Nebo Assignment" << ":\n";
1412  msg << "Left-hand side of assignment allocated ";
1413  msg << "on ";
1414  msg << "CPU but right-hand side is not ";
1415  msg << "(completely) accessible on the same CPU";
1416  msg << "\n";
1417  msg << "\t - " << __FILE__ << " : " << __LINE__;
1418  throw(std::runtime_error(msg.str()));
1419  }
1420  }
1421  else {
1422  std::ostringstream msg;
1423  msg << "Nebo error in " << "Nebo Assignment" << ":\n";
1424  msg << "Left-hand side of assignment allocated ";
1425  msg << "on ";
1426  msg << "unknown device - not on CPU or GPU";
1427  msg << "\n";
1428  msg << "\t - " << __FILE__ << " : " << __LINE__;
1429  throw(std::runtime_error(msg.str()));
1430  }
1431  }
1432 #endif
1433  }
1434 
1435 } // namespace SpatialOps
1436 #endif // NEBO_STENCIL_BUILDER_H
void operator()(SpatialMask< GammaFieldType > mask, PhiFieldType &phi, const GammaFieldType &gamma, bool minus) const
Apply boundary condition with gamma as a field.
NeboStencilBuilder(const CoefCollection &coefs)
construct a stencil with the specified coefficients
Provides tools to index into a sub-block of memory.
Definition: MemoryWindow.h:63
Supports definition of new Nebo stencils that do NOT invalidate ghost cells.
void apply_imp(MaskType mask, PhiFieldType &phi, const IntVec &shift, const NeboExpression< ExprType, GammaFieldType > &gamma, BCSide side) const
Apply boundary condition with gamma as an expression and a converted/non-converted mask...
CoefCollection const & coefs(void) const
Return coefficient collection.
OperatorT OperatorType
operator type
MaskShiftPoints< OperatorType, SrcFieldType, DestFieldType >::PlusPoint PlusPoint
positive face shift for mask
MinusResult minus(const SrcMask &src) const
Compute the minus side shift.
SrcFieldT SrcFieldType
source field type
Perform compile-time subtraction over a list of IndexTriplet types.
Definition: IndexTriplet.h:314
PntCltnT PointCollectionType
collection of stencil points
NeboBooleanExpression< PlusShift, DestFieldType > PlusResult
result type for positive shift
void operator()(SpatialMask< GammaFieldType > mask, PhiFieldType &phi, const NeboExpression< ExprType, GammaFieldType > &gamma, bool minus) const
Apply boundary condition with gamma as an expression.
PointCollection::AllButLast NonLowPoints
stencil offsets for all but low point in phi, assumes gamma is origin
Used for specifying field type traits.
Definition: IndexTriplet.h:121
void apply_to_field(const SrcFieldType &src, DestFieldType &dest) const
Apply this operator to the supplied source field to produce the supplied destination field...
Perform compile-time compare of two IndexTriplet types.
Definition: IndexTriplet.h:377
PntCltnT PointCollectionType
collection of stencil points
SrcFieldT SrcFieldType
source field type
Type for a Field Mask after it has been converted.
OperatorType type
operator type (Interpolant, Gradient, Divergence)
void apply_to_field(const SrcFieldType &src, DestFieldType &dest) const
Apply this operator to the supplied source field to produce the supplied destination field...
PlusResult plus(const SrcMask &src) const
Compute the plus side shift.
Defines a type to represent no direction.
OperatorType::PointCollectionType PointCollection
stencil point collection
NeboAverageStencilBuilder()
construct a stencil
NeboStencilCoefCollection< PointCollection::length > CoefCollection
collection of coefficients
NeboMaskShift< Initial, MinusPoint, Mask, DestFieldType > MinusShift
shift type for negative shift
Supports definition of new Nebo sum stencils, which sums given stencil points WITHOUT coefficients...
DestFieldT DestFieldType
destination field type
void apply_to_field(const SrcFieldType &src, DestFieldType &dest) const
Apply this operator to the supplied source field to produce the supplied destination field...
PointCollection::First HighPoint
stencil offset for high point in phi, assumes gamma is origin
SpatialMask< DestFieldType > DestMask
destination mask type
NeboBoundaryConditionBuilder(OperatorType const &op)
construct a boundary condition
NeboEdgelessStencilBuilder(const CoefCollection &coefs)
construct a stencil with the specified coefficients
for wide stencils where we set on a point rather than a face
SrcFieldT SrcFieldType
source field type
NeboMask< Initial, SrcFieldType > Mask
Nebo mask type.
void operator()(ConvertedPhiMask mask, PhiFieldType &phi, const NeboExpression< ExprType, GammaFieldType > &gamma) const
Apply boundary condition with gamma as an expression and a converted mask as boundary mask...
DestFieldT DestFieldType
destination field type
NeboStencilCoefCollection< PointCollectionType::length > CoefCollection
collection of coefficients
DestFieldT DestFieldType
destination field type
NeboSumStencilBuilder()
construct a stencil
OperatorType type
operator type (Interpolant, Gradient, Divergence)
Supports definition of new Nebo stencils.
Abstracts a mask.
Definition: SpatialMask.h:70
Perform compile-time compare of two IndexTriplet types.
Definition: IndexTriplet.h:356
OperatorType::DestFieldType GammaFieldType
type of fields in gamma, which this operator reads
NeboBooleanExpression< MinusShift, DestFieldType > MinusResult
result type for negative shift
PointCollection::AllButFirst NonHighPoints
stencil offsets for all but high point in phi, assumes gamma is origin
OperatorType::SrcFieldType PhiFieldType
field type of phi, which this operator modifies
IntVec get_plus() const
obtain the IntVec containing the number of ghost cells on the (+) faces
Definition: GhostData.h:145
PntCltnT PointCollectionType
collection of stencil points
NeboMaskShift< Initial, PlusPoint, Mask, DestFieldType > PlusShift
shift type for positive shift
void operator()(ConvertedPhiMask mask, PhiFieldType &phi, const double gamma) const
Apply boundary condition with gamma as a scalar.
Holds information about the number of ghost cells on each side of the domain.
Definition: GhostData.h:54
The Nebo Operator for switch.
Supports definition of new Nebo mask shift stencils, which converts a mask of one field type into ano...
Supports definition of new Nebo boundary condition.
SpatialMask< SrcFieldType > SrcMask
source mask type
SrcFieldT SrcFieldType
source field type
BCSide
Allows identification of whether we are setting the BC on the right or left side when using an operat...
NeboMaskShiftBuilder()
construct a stencil
PntCltnT PointCollectionType
collection of stencil points
PointCollection::Last LowPoint
stencil offset for low point in phi, assumes gamma is origin
void operator()(const SpatialMask< GammaFieldType > mask, PhiFieldType &phi, const double gamma, bool minus) const
Apply boundary condition with gamma as a scalar.
Supports definition of new Nebo average stencils, which automatically averages given stencil points...
NeboStencilCoefCollection< PointCollectionType::length > CoefCollection
collection of coefficients
SrcFieldT SrcFieldType
source field type
Abstracts a field.
Definition: SpatialField.h:132
void apply_to_field(const SrcFieldType &src, DestFieldType &dest) const
Apply this operator to the supplied source field to produce the supplied destination field...
IntVec get_minus() const
obtain the IntVec containing the number of ghost cells on the (-) faces
Definition: GhostData.h:135
void operator()(ConvertedPhiMask mask, PhiFieldType &phi, const GammaFieldType &gamma) const
Apply boundary condition with gamma as a field.
DestFieldT DestFieldType
destination field type
OperatorT OperatorType
type of operator to invert
CoefCollection const & coefs(void) const
Return coefficient collection.
MaskShiftPoints< OperatorType, SrcFieldType, DestFieldType >::MinusPoint MinusPoint
negative face shift for mask
DestFieldT DestFieldType
destination field type