SpatialOps
NeboStencils.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_STENCILS_H
26  #define NEBO_STENCILS_H
27 
28  namespace SpatialOps {
29  template<typename T, typename CollectionType>
31  public:
32  T typedef CurrentType;
33 
34  CollectionType typedef Collection;
35 
37 
38  CurrentType typedef First;
39 
40  Collection typedef AllButFirst;
41 
42  typename Collection::Last typedef Last;
43 
45  AllButLast;
46 
47  enum {length = 1 + Collection::length};
48 
49  template<typename NewType>
50  struct AddType {
52  };
53 
54  NeboGenericTypeList(Collection const & o, CurrentType const & c)
55  : others_(o), cur_(c)
56  {}
57 
58  template<typename NewType>
59  inline NeboGenericTypeList<NewType, MyType> const operator ()(NewType
60  const &
61  c) const {
62  return NeboGenericTypeList<NewType, MyType>(*this, c);
63  }
64 
65  #ifdef __CUDACC__
66  __host__ __device__
67  #endif
68  /* __CUDACC__ */ inline CurrentType current(void) const {
69  return cur_;
70  }
71 
72  #ifdef __CUDACC__
73  __host__ __device__
74  #endif
75  /* __CUDACC__ */ inline Collection const & others(void) const {
76  return others_;
77  }
78 
79  inline Last last(void) const { return others().last(); }
80 
81  inline AllButLast all_but_last(void) const {
82  return Collection::AllButLast(others().all_but_last(), cur_);
83  }
84 
85  private:
86  Collection const others_;
87 
88  CurrentType const cur_;
89  };
90 
92  public:
93  NeboGenericEmptyTypeList typedef MyType;
94 
95  enum {length = 0};
96 
97  template<typename NewType>
98  struct AddType {
100  };
101 
102  template<typename NewType>
103  inline NeboGenericTypeList<NewType, MyType> const operator ()(NewType
104  const &
105  c) const {
106  return NeboGenericTypeList<NewType, MyType>(*this, c);
107  }
108  };
109 
110  template<typename T>
112  public:
113  T typedef CurrentType;
114 
115  NeboGenericEmptyTypeList typedef Collection;
116 
118 
119  CurrentType typedef First;
120 
121  Collection typedef AllButFirst;
122 
123  First typedef Last;
124 
125  NeboGenericEmptyTypeList typedef AllButLast;
126 
127  enum {length = 1};
128 
129  template<typename NewType>
130  struct AddType {
132  };
133 
135  CurrentType const & c)
136  : others_(o), cur_(c)
137  {}
138 
139  template<typename NewType>
140  inline NeboGenericTypeList<NewType, MyType> const operator ()(NewType
141  const &
142  c) const {
143  return NeboGenericTypeList<NewType, MyType>(*this, c);
144  }
145 
146  #ifdef __CUDACC__
147  __host__ __device__
148  #endif
149  /* __CUDACC__ */ inline CurrentType current(void) const {
150  return cur_;
151  }
152 
153  #ifdef __CUDACC__
154  __host__ __device__
155  #endif
156  /* __CUDACC__ */ inline NeboGenericEmptyTypeList const & others(void) const {
157  return others_;
158  }
159 
160  inline Last last(void) const { return cur_; }
161 
162  inline AllButLast all_but_last(void) const { return others(); }
163 
164  private:
165  NeboGenericEmptyTypeList const others_;
166 
167  CurrentType const cur_;
168  };
169 
170  template<int Length>
172  public:
174  & o,
175  double const c)
176  : others_(o), coef_(c)
177  {}
178 
179  inline NeboStencilCoefCollection<Length + 1> const operator ()(double
180  const c) const {
181  return NeboStencilCoefCollection<Length + 1>(*this, c);
182  }
183 
184  #ifdef __CUDACC__
185  __host__ __device__
186  #endif
187  /* __CUDACC__ */ inline double coef(void) const { return coef_; }
188 
189  inline double get_coef(int const index) const {
190  if(index < 0) {
191  std::ostringstream msg;
192  msg << "Nebo error in " << "Nebo Stencil Coefficient Collection"
193  << ":\n";
194  msg << "given negative value for coefficient index";
195  msg << "\n";
196  msg << "\t - " << __FILE__ << " : " << __LINE__;
197  throw(std::runtime_error(msg.str()));
198  };
199 
200  if(index >= Length) {
201  std::ostringstream msg;
202  msg << "Nebo error in " << "Nebo Stencil Coefficient Collection"
203  << ":\n";
204  msg << "trying to access a coefficient that does not exist";
205  msg << "\n";
206  msg << "\t - " << __FILE__ << " : " << __LINE__;
207  throw(std::runtime_error(msg.str()));
208  };
209 
210  return (index == Length - 1 ? coef() : others().get_coef(index));
211  }
212 
213  #ifdef __CUDACC__
214  __host__ __device__
215  #endif
216  /* __CUDACC__ */ inline NeboStencilCoefCollection<Length - 1> const
217  others(void) const { return others_; }
218 
219  inline double last(void) const { return others().last(); }
220 
221  inline NeboStencilCoefCollection<Length - 1> all_but_last(void) const {
222  return NeboStencilCoefCollection<Length - 1>(others().all_but_last(),
223  coef_);
224  }
225 
226  private:
227  NeboStencilCoefCollection<Length - 1> const others_;
228 
229  double const coef_;
230  };
231 
232  template<>
234  public:
235  NeboStencilCoefCollection(double const c)
236  : coef_(c)
237  {}
238 
239  NeboStencilCoefCollection(NeboNil nil, double const c)
240  : coef_(c)
241  {}
242 
243  inline NeboStencilCoefCollection<2> const operator ()(double const c) const {
244  return NeboStencilCoefCollection<2>(*this, c);
245  }
246 
247  #ifdef __CUDACC__
248  __host__ __device__
249  #endif
250  /* __CUDACC__ */ inline double coef(void) const { return coef_; }
251 
252  inline double get_coef(int const index) const {
253  if(index < 0) {
254  std::ostringstream msg;
255  msg << "Nebo error in " << "Nebo Stencil Coefficient Collection"
256  << ":\n";
257  msg << "given negative value for coefficient index";
258  msg << "\n";
259  msg << "\t - " << __FILE__ << " : " << __LINE__;
260  throw(std::runtime_error(msg.str()));
261  };
262 
263  if(index > 1) {
264  std::ostringstream msg;
265  msg << "Nebo error in " << "Nebo Stencil Coefficient Collection"
266  << ":\n";
267  msg << "trying to access a coefficient that does not exist";
268  msg << "\n";
269  msg << "\t - " << __FILE__ << " : " << __LINE__;
270  throw(std::runtime_error(msg.str()));
271  };
272 
273  return coef();
274  }
275 
276  inline double last(void) const { return coef_; }
277 
278  inline NeboNil all_but_last(void) const { return NeboNil(); }
279 
280  private:
281  double const coef_;
282  };
283 
284  inline NeboStencilCoefCollection<1> const build_coef_collection(double
285  const c) {
287  };
288 
289  inline NeboStencilCoefCollection<2> const build_two_point_coef_collection(double
290  const
291  c1,
292  double
293  const
294  c2) {
295  return NeboStencilCoefCollection<1>(c1)(c2);
296  };
297 
298  inline NeboStencilCoefCollection<3> const
299  build_three_point_coef_collection(double const c1,
300  double const c2,
301  double const c3) {
302  return NeboStencilCoefCollection<1>(c1)(c2)(c3);
303  };
304 
305  inline NeboStencilCoefCollection<4> const build_four_point_coef_collection(double
306  const
307  c1,
308  double
309  const
310  c2,
311  double
312  const
313  c3,
314  double
315  const
316  c4) {
317  return NeboStencilCoefCollection<1>(c1)(c2)(c3)(c4);
318  };
319 
320  inline NeboStencilCoefCollection<5> const build_five_point_coef_collection(double
321  const
322  c1,
323  double
324  const
325  c2,
326  double
327  const
328  c3,
329  double
330  const
331  c4,
332  double
333  const
334  c5) {
335  return NeboStencilCoefCollection<1>(c1)(c2)(c3)(c4)(c5);
336  };
337 
338  inline NeboStencilCoefCollection<7> const
339  build_seven_point_coef_collection(double const c1,
340  double const c2,
341  double const c3,
342  double const c4,
343  double const c5,
344  double const c6,
345  double const c7) {
346  return NeboStencilCoefCollection<1>(c1)(c2)(c3)(c4)(c5)(c6)(c7);
347  };
348 
349  template<typename PointType, typename CollectionType>
351  public:
352  PointType typedef Point;
353 
354  CollectionType typedef Collection;
355 
357 
358  Point typedef First;
359 
360  Collection typedef AllButFirst;
361 
362  typename Collection::Last typedef Last;
363 
365  typedef AllButLast;
366 
367  enum {length = 1 + Collection::length};
368 
369  template<typename NewPoint>
370  struct AddPoint {
372  };
373 
374  static inline GhostData possible_ghosts(void) {
375  return min(additive_reductive_point_to_ghost(Point::int_vec()),
376  Collection::possible_ghosts());
377  }
378 
379  static inline GhostData possible_ghosts(GhostData const & ghosts) {
380  return ghosts + possible_ghosts();
381  }
382 
383  static inline GhostData possible_additive_ghosts(void) {
384  return min(addative_point_to_ghost(Point::int_vec()),
385  Collection::possible_ghosts());
386  }
387 
388  static inline GhostData possible_additive_ghosts(GhostData const &
389  ghosts) {
390  return ghosts + possible_additive_ghosts();
391  }
392  };
393 
394  template<typename PointType>
395  struct NeboStencilPointCollection<PointType, NeboNil> {
396  public:
397  PointType typedef Point;
398 
399  NeboNil typedef Collection;
400 
402 
403  Point typedef Last;
404 
405  NeboNil typedef AllButLast;
406 
407  enum {length = 1};
408 
409  template<typename NewPoint>
410  struct AddPoint {
412  };
413 
414  static inline GhostData possible_ghosts(void) {
415  return additive_reductive_point_to_ghost(Point::int_vec());
416  }
417 
418  static inline GhostData possible_ghosts(GhostData const & ghosts) {
419  return ghosts + possible_ghosts();
420  }
421 
422  static inline GhostData possible_additive_ghosts(void) {
423  return additive_point_to_ghost(Point::int_vec());
424  }
425 
426  static inline GhostData possible_additive_ghosts(GhostData const &
427  ghosts) {
428  return ghosts + possible_additive_ghosts();
429  }
430  };
431 
432  template<typename CurrentMode,
433  typename Pts,
434  typename Arg,
435  typename FieldType>
436  struct NeboStencil;
437  template<typename Pts, typename Arg, typename FieldType>
438  struct NeboStencil<Initial, Pts, Arg, FieldType> {
439  public:
440  FieldType typedef field_type;
441 
443 
445  typedef SeqWalkType;
446 
447  #ifdef ENABLE_THREADS
449  typedef ResizeType;
450  #endif
451  /* ENABLE_THREADS */
452 
453  #ifdef __CUDACC__
455  typedef GPUWalkType;
456  #endif
457  /* __CUDACC__ */
458 
459  NeboStencil(Arg const & a, Coefs const & coefs)
460  : arg_(a), coefs_(coefs)
461  {}
462 
463  inline GhostData ghosts_with_bc(void) const {
464  return Pts::possible_ghosts(arg_.ghosts_with_bc());
465  }
466 
467  inline GhostData ghosts_without_bc(void) const {
468  return Pts::possible_ghosts(arg_.ghosts_without_bc());
469  }
470 
471  inline bool has_extents(void) const { return arg_.has_extents(); }
472 
473  inline IntVec extents(void) const { return arg_.extents(); }
474 
475  inline IntVec has_bc(BCSide const bcSide) const {
476  return arg_.has_bc(bcSide);
477  }
478 
479  inline SeqWalkType init(IntVec const & extents,
480  GhostData const & ghosts,
481  IntVec const & hasBC) const {
482  return SeqWalkType(arg_.init(extents, ghosts, hasBC), coefs_);
483  }
484 
485  #ifdef ENABLE_THREADS
486  inline ResizeType resize(void) const {
487  return ResizeType(arg_.resize(), coefs_);
488  }
489  #endif
490  /* ENABLE_THREADS */
491 
492  #ifdef __CUDACC__
493  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
494 
495  inline bool gpu_ready(int const deviceIndex) const {
496  return arg_.gpu_ready(deviceIndex);
497  }
498 
499  inline GPUWalkType gpu_init(IntVec const & extents,
500  GhostData const & ghosts,
501  IntVec const & hasBC,
502  int const deviceIndex,
503  cudaStream_t const & lhsStream) const {
504  return GPUWalkType(arg_.gpu_init(extents,
505  ghosts,
506  hasBC,
507  deviceIndex,
508  lhsStream),
509  coefs_);
510  }
511 
512  inline void stream_wait_event(cudaEvent_t const & event) const {
513  arg_.stream_wait_event(event);
514  }
515 
516  #ifdef NEBO_GPU_TEST
517  inline void gpu_prep(int const deviceIndex) const {
518  arg_.gpu_prep(deviceIndex);
519  }
520  #endif
521  /* NEBO_GPU_TEST */
522  #endif
523  /* __CUDACC__ */
524 
525  private:
526  Arg const arg_;
527 
528  Coefs const coefs_;
529  };
530  #ifdef ENABLE_THREADS
531  template<typename Pts, typename Arg, typename FieldType>
532  struct NeboStencil<Resize, Pts, Arg, FieldType> {
533  public:
534  FieldType typedef field_type;
535 
537 
539  typedef SeqWalkType;
540 
541  NeboStencil(Arg const & arg, Coefs const & coefs)
542  : arg_(arg), coefs_(coefs)
543  {}
544 
545  inline SeqWalkType init(IntVec const & extents,
546  GhostData const & ghosts,
547  IntVec const & hasBC) const {
548  return SeqWalkType(arg_.init(extents, ghosts, hasBC), coefs_);
549  }
550 
551  private:
552  Arg const arg_;
553 
554  Coefs const coefs_;
555  }
556  #endif
557  /* ENABLE_THREADS */;
558  template<typename Pts, typename Arg, typename FieldType>
559  struct NeboStencil<SeqWalk, Pts, Arg, FieldType> {
560  public:
561  FieldType typedef field_type;
562 
563  typename field_type::value_type typedef value_type;
564 
566 
567  template<typename PointCollection>
568  struct EvalExpr {
570 
571  typename PointCollection::Point typedef Point;
572 
573  typename PointCollection::Collection typedef Collection;
574 
575  template<typename OptionalArgT>
576  static inline value_type eval(Arg const & arg,
577  Coefs const & coefs,
578  int const x,
579  int const y,
580  int const z) {
581  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
582  coefs.others(),
583  x,
584  y,
585  z) +
586  arg.template eval<OptionalArgT>(x + Point::value(0),
587  y + Point::value(1),
588  z + Point::value(2)) *
589  coefs.coef();
590  }
591  };
592 
593  template<typename Point>
594  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
595  NeboStencilCoefCollection<1> typedef Coefs;
596 
597  template<typename OptionalArgT>
598  static inline value_type eval(Arg const & arg,
599  Coefs const & coefs,
600  int const x,
601  int const y,
602  int const z) {
603  return arg.template eval<OptionalArgT>(x + Point::value(0),
604  y + Point::value(1),
605  z + Point::value(2)) *
606  coefs.coef();
607  }
608  };
609 
610  NeboStencil(Arg const & arg, Coefs const & coefs)
611  : arg_(arg), coefs_(coefs)
612  {}
613 
614  template<typename OptionalArgT>
615  inline value_type eval(int const x, int const y, int const z) const {
616  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
617  coefs_,
618  x,
619  y,
620  z);
621  }
622 
623  private:
624  Arg arg_;
625 
626  Coefs const coefs_;
627  };
628  #ifdef __CUDACC__
629  template<typename Pts, typename Arg, typename FieldType>
630  struct NeboStencil<GPUWalk, Pts, Arg, FieldType> {
631  public:
632  FieldType typedef field_type;
633 
634  typename field_type::value_type typedef value_type;
635 
637 
638  template<typename PointCollection>
639  struct EvalExpr {
641  ;
642 
643  typename PointCollection::Point typedef Point;
644 
645  typename PointCollection::Collection typedef Collection;
646 
647  template<typename OptionalArgT>
648  __device__ static inline value_type eval(Arg const & arg,
649  Coefs const & coefs,
650  int const x,
651  int const y,
652  int const z) {
653  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
654  coefs.others(),
655  x,
656  y,
657  z)
658  + arg.template eval<OptionalArgT>(x + Point::
659  value_gpu(0),
660  y + Point::
661  value_gpu(1),
662  z + Point::
663  value_gpu(2)) *
664  coefs.coef();
665  }
666  };
667 
668  template<typename Point>
669  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
670  NeboStencilCoefCollection<1> typedef Coefs;
671 
672  template<typename OptionalArgT>
673  __device__ static inline value_type eval(Arg const & arg,
674  Coefs const & coefs,
675  int const x,
676  int const y,
677  int const z) {
678  return arg.template eval<OptionalArgT>(x + Point::value_gpu(0),
679  y + Point::value_gpu(1),
680  z + Point::value_gpu(2))
681  * coefs.coef();
682  }
683  };
684 
685  NeboStencil(Arg const & a, Coefs const & coefs)
686  : arg_(a), coefs_(coefs)
687  {}
688 
689  template<typename OptionalArgT>
690  __device__ inline value_type eval(int const x,
691  int const y,
692  int const z) const {
693  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
694  coefs_,
695  x,
696  y,
697  z);
698  }
699 
700  private:
701  Arg arg_;
702 
703  Coefs const coefs_;
704  }
705  #endif
706  /* __CUDACC__ */;
707 
708  template<typename CurrentMode,
709  typename Pts,
710  typename Arg,
711  typename FieldType>
713  template<typename Pts, typename Arg, typename FieldType>
714  struct NeboEdgelessStencil<Initial, Pts, Arg, FieldType> {
715  public:
716  FieldType typedef field_type;
717 
719 
721  typedef SeqWalkType;
722 
723  #ifdef ENABLE_THREADS
724  NeboEdgelessStencil<Resize,
725  Pts,
726  typename Arg::ResizeType,
727  FieldType> typedef ResizeType;
728  #endif
729  /* ENABLE_THREADS */
730 
731  #ifdef __CUDACC__
732  NeboEdgelessStencil<GPUWalk,
733  Pts,
734  typename Arg::GPUWalkType,
735  FieldType> typedef GPUWalkType;
736  #endif
737  /* __CUDACC__ */
738 
739  NeboEdgelessStencil(Arg const & a, Coefs const & coefs)
740  : arg_(a), coefs_(coefs)
741  {}
742 
743  inline GhostData ghosts_with_bc(void) const {
744  return Pts::possible_additive_ghosts(arg_.ghosts_with_bc());
745  }
746 
747  inline GhostData ghosts_without_bc(void) const {
748  return Pts::possible_additive_ghosts(arg_.ghosts_without_bc());
749  }
750 
751  inline bool has_extents(void) const { return arg_.has_extent(); }
752 
753  inline IntVec extents(void) const { return arg_.extents(); }
754 
755  inline IntVec has_bc(BCSide const bcSide) const {
756  return arg_.has_bc(bcSide);
757  }
758 
759  inline SeqWalkType init(IntVec const & extents,
760  GhostData const & ghosts,
761  IntVec const & hasBC) const {
762  return SeqWalkType(arg_.init(extents, ghosts, hasBC),
763  coefs_,
764  lowest_indicies(),
765  highest_indicies());
766  }
767 
768  #ifdef ENABLE_THREADS
769  inline ResizeType resize(void) const {
770  return ResizeType(arg_.resize(),
771  coefs_,
772  lowest_indicies(),
773  highest_indicies());
774  }
775  #endif
776  /* ENABLE_THREADS */
777 
778  #ifdef __CUDACC__
779  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
780 
781  inline bool gpu_ready(int const deviceIndex) const {
782  return arg_.gpu_ready(deviceIndex);
783  }
784 
785  inline GPUWalkType gpu_init(IntVec const & extents,
786  GhostData const & ghosts,
787  IntVec const & hasBC,
788  int const deviceIndex,
789  cudaStream_t const & lhsStream) const {
790  return GPUWalkType(arg_.gpu_init(extents,
791  ghosts,
792  hasBC,
793  deviceIndex,
794  lhsStream),
795  coefs_);
796  }
797 
798  inline void stream_wait_event(cudaEvent_t const & event) const {
799  arg_.stream_wait_event(event);
800  }
801 
802  #ifdef NEBO_GPU_TEST
803  inline void gpu_prep(int const deviceIndex) const {
804  arg_.gpu_prep(deviceIndex);
805  }
806  #endif
807  /* NEBO_GPU_TEST */
808  #endif
809  /* __CUDACC__ */
810 
811  inline GhostData actual_ghosts(void) const {
812  return Pts::possible_ghosts(arg_.ghosts_with_bc());
813  }
814 
815  inline IntVec lowest_indicies(void) const {
816  return -(actual_ghosts().get_minus());
817  }
818 
819  inline IntVec highest_indicies(void) const {
820  return actual_ghosts().get_plus() + extents();
821  }
822 
823  private:
824  Arg const arg_;
825 
826  Coefs const coefs_;
827  };
828  #ifdef ENABLE_THREADS
829  template<typename Pts, typename Arg, typename FieldType>
830  struct NeboEdgelessStencil<Resize, Pts, Arg, FieldType> {
831  public:
832  FieldType typedef field_type;
833 
835 
836  NeboEdgelessStencil<SeqWalk,
837  Pts,
838  typename Arg::SeqWalkType,
839  FieldType> typedef SeqWalkType;
840 
841  NeboEdgelessStencil(Arg const & arg,
842  Coefs const & coefs,
843  IntVec const & low,
844  IntVec const & high)
845  : arg_(arg), coefs_(coefs), low_(low), high_(high)
846  {}
847 
848  inline SeqWalkType init(IntVec const & extents,
849  GhostData const & ghosts,
850  IntVec const & hasBC) const {
851  return SeqWalkType(arg_.init(extents, ghosts, hasBC),
852  coefs_,
853  low_,
854  high_);
855  }
856 
857  private:
858  Arg const arg_;
859 
860  Coefs const coefs_;
861 
862  IntVec const low_;
863 
864  IntVec const high_;
865  }
866  #endif
867  /* ENABLE_THREADS */;
868  template<typename Pts, typename Arg, typename FieldType>
869  struct NeboEdgelessStencil<SeqWalk, Pts, Arg, FieldType> {
870  public:
871  FieldType typedef field_type;
872 
873  typename field_type::value_type typedef value_type;
874 
876 
877  template<typename PointCollection>
878  struct EvalExpr {
880 
881  typename PointCollection::Point typedef Point;
882 
883  typename PointCollection::Collection typedef Collection;
884 
885  template<typename OptionalArgT>
886  static inline value_type eval(Arg const & arg,
887  Coefs const & coefs,
888  int const x,
889  int const y,
890  int const z) {
891  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
892  coefs.others(),
893  x,
894  y,
895  z) +
896  arg.template eval<OptionalArgT>(x + Point::value(0),
897  y + Point::value(1),
898  z + Point::value(2)) *
899  coefs.coef();
900  }
901  };
902 
903  template<typename Point>
904  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
905  NeboStencilCoefCollection<1> typedef Coefs;
906 
907  template<typename OptionalArgT>
908  static inline value_type eval(Arg const & arg,
909  Coefs const & coefs,
910  int const x,
911  int const y,
912  int const z) {
913  return arg.template eval<OptionalArgT>(x + Point::value(0),
914  y + Point::value(1),
915  z + Point::value(2)) *
916  coefs.coef();
917  }
918  };
919 
920  NeboEdgelessStencil(Arg const & arg,
921  Coefs const & coefs,
922  IntVec const & low,
923  IntVec const & high)
924  : arg_(arg), coefs_(coefs), low_(low), high_(high)
925  {}
926 
927  template<typename OptionalArgT>
928  inline value_type eval(int const x, int const y, int const z) const {
929  #ifndef NDEBUG
930  IntVec index = IntVec(x, y, z);
931  if(index < low_ || index >= high_) {
932  std::ostringstream msg;
933  msg << "Nebo error in " << "Nebo Edgeless Stencil" << ":\n";
934  msg << " - " << low_ << " < " << index << " <= " << high_;
935  msg << "\n";
936  msg << "\t - " << __FILE__ << " : " << __LINE__;
937  throw(std::runtime_error(msg.str()));;
938  }
939  #endif
940  /* NDEBUG */;
941 
942  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
943  coefs_,
944  x,
945  y,
946  z);
947  }
948 
949  private:
950  Arg arg_;
951 
952  Coefs const coefs_;
953 
954  IntVec const low_;
955 
956  IntVec const high_;
957  };
958  #ifdef __CUDACC__
959  template<typename Pts, typename Arg, typename FieldType>
960  struct NeboEdgelessStencil<GPUWalk, Pts, Arg, FieldType> {
961  public:
962  FieldType typedef field_type;
963 
964  typename field_type::value_type typedef value_type;
965 
967 
968  template<typename PointCollection>
969  struct EvalExpr {
971  ;
972 
973  typename PointCollection::Point typedef Point;
974 
975  typename PointCollection::Collection typedef Collection;
976 
977  template<typename OptionalArgT>
978  __device__ static inline value_type eval(Arg const & arg,
979  Coefs const & coefs,
980  int const x,
981  int const y,
982  int const z) {
983  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
984  coefs.others(),
985  x,
986  y,
987  z)
988  + arg.template eval<OptionalArgT>(x + Point::
989  value_gpu(0),
990  y + Point::
991  value_gpu(1),
992  z + Point::
993  value_gpu(2)) *
994  coefs.coef();
995  }
996  };
997 
998  template<typename Point>
999  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
1000  NeboStencilCoefCollection<1> typedef Coefs;
1001 
1002  template<typename OptionalArgT>
1003  __device__ static inline value_type eval(Arg const & arg,
1004  Coefs const & coefs,
1005  int const x,
1006  int const y,
1007  int const z) {
1008  return arg.template eval<OptionalArgT>(x + Point::value_gpu(0),
1009  y + Point::value_gpu(1),
1010  z + Point::value_gpu(2))
1011  * coefs.coef();
1012  }
1013  };
1014 
1015  NeboEdgelessStencil(Arg const & a, Coefs const & coefs)
1016  : arg_(a), coefs_(coefs)
1017  {}
1018 
1019  template<typename OptionalArgT>
1020  __device__ inline value_type eval(int const x,
1021  int const y,
1022  int const z) const {
1023  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
1024  coefs_,
1025  x,
1026  y,
1027  z);
1028  }
1029 
1030  private:
1031  Arg arg_;
1032 
1033  Coefs const coefs_;
1034  }
1035  #endif
1036  /* __CUDACC__ */;
1037 
1038  template<typename CurrentMode,
1039  typename Pts,
1040  typename Arg,
1041  typename FieldType>
1043  template<typename Pts, typename Arg, typename FieldType>
1044  struct NeboSumStencil<Initial, Pts, Arg, FieldType> {
1045  public:
1046  FieldType typedef field_type;
1047 
1049  typedef SeqWalkType;
1050 
1051  #ifdef ENABLE_THREADS
1053  typedef ResizeType;
1054  #endif
1055  /* ENABLE_THREADS */
1056 
1057  #ifdef __CUDACC__
1059  typedef GPUWalkType;
1060  #endif
1061  /* __CUDACC__ */
1062 
1063  NeboSumStencil(Arg const & a)
1064  : arg_(a)
1065  {}
1066 
1067  inline GhostData ghosts_with_bc(void) const {
1068  return Pts::possible_ghosts(arg_.ghosts_with_bc());
1069  }
1070 
1071  inline GhostData ghosts_without_bc(void) const {
1072  return Pts::possible_ghosts(arg_.ghosts_without_bc());
1073  }
1074 
1075  inline bool has_extents(void) const { return arg_.has_extents(); }
1076 
1077  inline IntVec extents(void) const { return arg_.extents(); }
1078 
1079  inline IntVec has_bc(BCSide const bcSide) const {
1080  return arg_.has_bc(bcSide);
1081  }
1082 
1083  inline SeqWalkType init(IntVec const & extents,
1084  GhostData const & ghosts,
1085  IntVec const & hasBC) const {
1086  return SeqWalkType(arg_.init(extents, ghosts, hasBC));
1087  }
1088 
1089  #ifdef ENABLE_THREADS
1090  inline ResizeType resize(void) const {
1091  return ResizeType(arg_.resize());
1092  }
1093  #endif
1094  /* ENABLE_THREADS */
1095 
1096  #ifdef __CUDACC__
1097  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
1098 
1099  inline bool gpu_ready(int const deviceIndex) const {
1100  return arg_.gpu_ready(deviceIndex);
1101  }
1102 
1103  inline GPUWalkType gpu_init(IntVec const & extents,
1104  GhostData const & ghosts,
1105  IntVec const & hasBC,
1106  int const deviceIndex,
1107  cudaStream_t const & lhsStream) const {
1108  return GPUWalkType(arg_.gpu_init(extents,
1109  ghosts,
1110  hasBC,
1111  deviceIndex,
1112  lhsStream));
1113  }
1114 
1115  inline void stream_wait_event(cudaEvent_t const & event) const {
1116  arg_.stream_wait_event(event);
1117  }
1118 
1119  #ifdef NEBO_GPU_TEST
1120  inline void gpu_prep(int const deviceIndex) const {
1121  arg_.gpu_prep(deviceIndex);
1122  }
1123  #endif
1124  /* NEBO_GPU_TEST */
1125  #endif
1126  /* __CUDACC__ */
1127 
1128  private:
1129  Arg const arg_;
1130  };
1131  #ifdef ENABLE_THREADS
1132  template<typename Pts, typename Arg, typename FieldType>
1133  struct NeboSumStencil<Resize, Pts, Arg, FieldType> {
1134  public:
1135  FieldType typedef field_type;
1136 
1138  typedef SeqWalkType;
1139 
1140  NeboSumStencil(Arg const & arg)
1141  : arg_(arg)
1142  {}
1143 
1144  inline SeqWalkType init(IntVec const & extents,
1145  GhostData const & ghosts,
1146  IntVec const & hasBC) const {
1147  return SeqWalkType(arg_.init(extents, ghosts, hasBC));
1148  }
1149 
1150  private:
1151  Arg const arg_;
1152  }
1153  #endif
1154  /* ENABLE_THREADS */;
1155  template<typename Pts, typename Arg, typename FieldType>
1156  struct NeboSumStencil<SeqWalk, Pts, Arg, FieldType> {
1157  public:
1158  FieldType typedef field_type;
1159 
1160  typename field_type::value_type typedef value_type;
1161 
1162  template<typename PointCollection>
1163  struct EvalExpr {
1164  typename PointCollection::Point typedef Point;
1165 
1166  typename PointCollection::Collection typedef Collection;
1167 
1168  template<typename OptionalArgT>
1169  static inline value_type eval(Arg const & arg,
1170  int const x,
1171  int const y,
1172  int const z) {
1173  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
1174  x,
1175  y,
1176  z) +
1177  arg.template eval<OptionalArgT>(x + Point::value(0),
1178  y + Point::value(1),
1179  z + Point::value(2));
1180  }
1181  };
1182 
1183  template<typename Point>
1184  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
1185  template<typename OptionalArgT>
1186  static inline value_type eval(Arg const & arg,
1187  int const x,
1188  int const y,
1189  int const z) {
1190  return arg.template eval<OptionalArgT>(x + Point::value(0),
1191  y + Point::value(1),
1192  z + Point::value(2));
1193  }
1194  };
1195 
1196  NeboSumStencil(Arg const & arg)
1197  : arg_(arg)
1198  {}
1199 
1200  template<typename OptionalArgT>
1201  inline value_type eval(int const x, int const y, int const z) const {
1202  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_, x, y, z);
1203  }
1204 
1205  private:
1206  Arg arg_;
1207  };
1208  #ifdef __CUDACC__
1209  template<typename Pts, typename Arg, typename FieldType>
1210  struct NeboSumStencil<GPUWalk, Pts, Arg, FieldType> {
1211  public:
1212  FieldType typedef field_type;
1213 
1214  typename field_type::value_type typedef value_type;
1215 
1216  template<typename PointCollection>
1217  struct EvalExpr {
1218  typename PointCollection::Point typedef Point;
1219 
1220  typename PointCollection::Collection typedef Collection;
1221 
1222  template<typename OptionalArgT>
1223  __device__ static inline value_type eval(Arg const & arg,
1224  int const x,
1225  int const y,
1226  int const z) {
1227  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
1228  x,
1229  y,
1230  z)
1231  + arg.template eval<OptionalArgT>(x + Point::
1232  value_gpu(0),
1233  y + Point::
1234  value_gpu(1),
1235  z + Point::
1236  value_gpu(2));
1237  }
1238  };
1239 
1240  template<typename Point>
1241  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
1242  template<typename OptionalArgT>
1243  __device__ static inline value_type eval(Arg const & arg,
1244  int const x,
1245  int const y,
1246  int const z) {
1247  return arg.template eval<OptionalArgT>(x + Point::value_gpu(0),
1248  y + Point::value_gpu(1),
1249  z + Point::value_gpu(2));
1250  }
1251  };
1252 
1253  NeboSumStencil(Arg const & a)
1254  : arg_(a)
1255  {}
1256 
1257  template<typename OptionalArgT>
1258  __device__ inline value_type eval(int const x,
1259  int const y,
1260  int const z) const {
1261  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_, x, y, z);
1262  }
1263 
1264  private:
1265  Arg arg_;
1266  }
1267  #endif
1268  /* __CUDACC__ */;
1269 
1270  template<typename Point>
1271  static inline GhostData point_possible_ghosts(GhostData const & ghosts) {
1272  return ghosts + additive_reductive_point_to_ghost(Point::int_vec());
1273  };
1274 
1275  template<typename CurrentMode,
1276  typename Point,
1277  typename Arg,
1278  typename FieldType>
1280  template<typename Point, typename Arg, typename FieldType>
1281  struct NeboMaskShift<Initial, Point, Arg, FieldType> {
1282  public:
1283  FieldType typedef field_type;
1284 
1285  typename Arg::SeqWalkType typedef ArgSeqWalkType;
1286 
1287  #ifdef __CUDACC__
1288  typename Arg::GPUWalkType typedef ArgGPUWalkType;
1289  #endif
1290  /* __CUDACC__ */
1291 
1293  SeqWalkType;
1294 
1295  #ifdef ENABLE_THREADS
1297  typedef ResizeType;
1298  #endif
1299  /* ENABLE_THREADS */
1300 
1301  #ifdef __CUDACC__
1303  GPUWalkType;
1304  #endif
1305  /* __CUDACC__ */
1306 
1307  NeboMaskShift(Arg const & a)
1308  : arg_(a)
1309  {}
1310 
1311  inline GhostData ghosts_with_bc(void) const {
1312  return point_possible_ghosts<Point>(arg_.ghosts_with_bc());
1313  }
1314 
1315  inline GhostData ghosts_without_bc(void) const {
1316  return point_possible_ghosts<Point>(arg_.ghosts_without_bc());
1317  }
1318 
1319  inline bool has_extents(void) const { return arg_.has_extents(); }
1320 
1321  inline IntVec extents(void) const { return arg_.extents(); }
1322 
1323  inline IntVec has_bc(BCSide const bcSide) const {
1324  return arg_.has_bc(bcSide);
1325  }
1326 
1327  inline SeqWalkType init(IntVec const & extents,
1328  GhostData const & ghosts,
1329  IntVec const & hasBC) const {
1330  return SeqWalkType(arg_.init(extents, ghosts, hasBC));
1331  }
1332 
1333  #ifdef ENABLE_THREADS
1334  inline ResizeType resize(void) const {
1335  return ResizeType(arg_.resize());
1336  }
1337  #endif
1338  /* ENABLE_THREADS */
1339 
1340  #ifdef __CUDACC__
1341  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
1342 
1343  inline bool gpu_ready(int const deviceIndex) const {
1344  return arg_.gpu_ready(deviceIndex);
1345  }
1346 
1347  inline GPUWalkType gpu_init(IntVec const & extents,
1348  GhostData const & ghosts,
1349  IntVec const & hasBC,
1350  int const deviceIndex,
1351  cudaStream_t const & lhsStream) const {
1352  return GPUWalkType(arg_.gpu_init(extents,
1353  ghosts,
1354  hasBC,
1355  deviceIndex,
1356  lhsStream));
1357  }
1358 
1359  inline void stream_wait_event(cudaEvent_t const & event) const {
1360  arg_.stream_wait_event(event);
1361  }
1362 
1363  #ifdef NEBO_GPU_TEST
1364  inline void gpu_prep(int const deviceIndex) const {
1365  arg_.gpu_prep(deviceIndex);
1366  }
1367  #endif
1368  /* NEBO_GPU_TEST */
1369  #endif
1370  /* __CUDACC__ */
1371 
1372  private:
1373  Arg const arg_;
1374  };
1375  #ifdef ENABLE_THREADS
1376  template<typename Point, typename Arg, typename FieldType>
1377  struct NeboMaskShift<Resize, Point, Arg, FieldType> {
1378  public:
1379  FieldType typedef field_type;
1380 
1381  typename Arg::SeqWalkType typedef ArgSeqWalkType;
1382 
1384  SeqWalkType;
1385 
1386  NeboMaskShift(Arg const & arg)
1387  : arg_(arg)
1388  {}
1389 
1390  inline SeqWalkType init(IntVec const & extents,
1391  GhostData const & ghosts,
1392  IntVec const & hasBC) const {
1393  return SeqWalkType(arg_.init(extents, ghosts, hasBC));
1394  }
1395 
1396  private:
1397  Arg const arg_;
1398  }
1399  #endif
1400  /* ENABLE_THREADS */;
1401  template<typename Point, typename Arg, typename FieldType>
1402  struct NeboMaskShift<SeqWalk, Point, Arg, FieldType> {
1403  public:
1404  FieldType typedef field_type;
1405 
1406  typename field_type::value_type typedef value_type;
1407 
1408  NeboMaskShift(Arg const & arg)
1409  : arg_(arg)
1410  {}
1411 
1412  template<typename OptionalArgT>
1413  inline bool eval(int const x, int const y, int const z) const {
1414  return arg_.template eval<OptionalArgT>(x + Point::value(0),
1415  y + Point::value(1),
1416  z + Point::value(2));
1417  }
1418 
1419  private:
1420  Arg arg_;
1421  };
1422  #ifdef __CUDACC__
1423  template<typename Point, typename Arg, typename FieldType>
1424  struct NeboMaskShift<GPUWalk, Point, Arg, FieldType> {
1425  public:
1426  FieldType typedef field_type;
1427 
1428  typename field_type::value_type typedef value_type;
1429 
1430  NeboMaskShift(Arg const & a)
1431  : arg_(a)
1432  {}
1433 
1434  template<typename OptionalArgT>
1435  __device__ inline bool eval(int const x, int const y, int const z) const {
1436  return arg_.template eval<OptionalArgT>(x + Point::value_gpu(0),
1437  y + Point::value_gpu(1),
1438  z + Point::value_gpu(2));
1439  }
1440 
1441  private:
1442  Arg arg_;
1443  }
1444  #endif
1445  /* __CUDACC__ */;
1446 
1447  template<typename CurrentMode,
1448  typename StencilT,
1449  typename NegativeStencilTs,
1450  typename PositiveStencilTs,
1451  typename Arg,
1452  typename DirT,
1453  typename FieldType>
1455  template<typename StencilT,
1456  typename NegativeStencilTs,
1457  typename PositiveStencilTs,
1458  typename Arg,
1459  typename DirT,
1460  typename FieldType>
1462  StencilT,
1463  NegativeStencilTs,
1464  PositiveStencilTs,
1465  Arg,
1466  DirT,
1467  FieldType> {
1468  public:
1469  FieldType typedef field_type;
1470 
1471  template<typename ListT, typename Dummy = void>
1472  struct ConvertToSeqWalk {
1473  public:
1474  NeboGenericTypeList<typename ListT::First::SeqWalkType,
1475  typename ConvertToSeqWalk<typename ListT::
1476  AllButFirst>::Result>
1477  typedef Result;
1478 
1479  static inline Result convert_runtime(ListT const list,
1480  IntVec const & extents,
1481  GhostData const & ghosts,
1482  IntVec const & hasBC) {
1483  return ConvertToSeqWalk<typename ListT::AllButFirst>::
1484  convert_runtime(list.others(), extents, ghosts, hasBC)(list.current().init(extents,
1485  ghosts,
1486  hasBC));
1487  }
1488  };
1489  template<typename Dummy>
1490  struct ConvertToSeqWalk<NeboGenericEmptyTypeList, Dummy> {
1491  public:
1492  NeboGenericEmptyTypeList typedef Result;
1493 
1494  static inline Result convert_runtime(NeboGenericEmptyTypeList
1495  const list,
1496  IntVec const & extents,
1497  GhostData const & ghosts,
1498  IntVec const & hasBC) {
1499  return list;
1500  }
1501  };
1502 
1503  template<typename ListT, typename Dummy = void>
1504  struct ConvertToResize {
1505  public:
1506  NeboGenericTypeList<typename ListT::First::ResizeType,
1507  typename ConvertToResize<typename ListT::
1508  AllButFirst>::Result>
1509  typedef Result;
1510 
1511  static inline Result convert_runtime(ListT const list) {
1512  return ConvertToResize<typename ListT::AllButFirst>::
1513  convert_runtime(list.others())(list.current().resize());
1514  }
1515  };
1516  template<typename Dummy>
1517  struct ConvertToResize<NeboGenericEmptyTypeList, Dummy> {
1518  public:
1519  NeboGenericEmptyTypeList typedef Result;
1520 
1521  static inline Result convert_runtime(NeboGenericEmptyTypeList
1522  const list) { return list; }
1523  };
1524 
1525  #ifdef __CUDACC__
1526  template<typename ListT, typename Dummy = void>
1527  struct ConvertToGPUWalk {
1528  public:
1529  NeboGenericTypeList<typename ListT::First::GPUWalkType,
1530  typename ConvertToGPUWalk<typename ListT::
1531  AllButFirst>::
1532  Result> typedef Result;
1533 
1534  static inline Result convert_runtime(ListT const list,
1535  IntVec const & extents,
1536  GhostData const & ghosts,
1537  IntVec const & hasBC,
1538  int const deviceIndex,
1539  cudaStream_t const &
1540  lhsStream) {
1541  return ConvertToGPUWalk<typename ListT::AllButFirst>::
1542  convert_runtime(list.others(),
1543  extents,
1544  ghosts,
1545  hasBC,
1546  deviceIndex,
1547  lhsStream)(list.current().gpu_init(extents,
1548  ghosts,
1549  hasBC,
1550  deviceIndex,
1551  lhsStream));
1552  }
1553  };
1554  template<typename Dummy>
1555  struct ConvertToGPUWalk<NeboGenericEmptyTypeList, Dummy> {
1556  public:
1557  NeboGenericEmptyTypeList typedef Result;
1558 
1559  static inline Result convert_runtime(NeboGenericEmptyTypeList
1560  const list,
1561  IntVec const & extents,
1562  GhostData const & ghosts,
1563  IntVec const & hasBC,
1564  int const deviceIndex,
1565  cudaStream_t const &
1566  lhsStream) { return list; }
1567  };
1568  #endif
1569  /* __CUDACC__ */
1570 
1571  template<typename StencilListT,
1572  typename MinusOrPlusT,
1573  typename Dummy = void>
1574  struct VerifyNoGhostUseOnEdge;
1575  template<typename StencilListT, typename Dummy>
1576  struct VerifyNoGhostUseOnEdge<StencilListT,
1578  Dummy> {
1579  static inline bool verify(StencilListT const & stencils,
1580  GhostData const & baseArgGhost,
1581  bool const previous) {
1582  return VerifyNoGhostUseOnEdge<typename StencilListT::Collection,
1584  verify(stencils.others(),
1585  baseArgGhost,
1586  previous && baseArgGhost.get_minus(DirT::value) -
1587  stencils.current().ghosts_with_bc().get_minus(DirT::
1588  value)
1589  < StencilListT::length);
1590  }
1591  };
1592  template<typename Dummy>
1593  struct VerifyNoGhostUseOnEdge<NeboGenericEmptyTypeList,
1595  Dummy> {
1596  static inline bool verify(NeboGenericEmptyTypeList const & stencils,
1597  GhostData const & baseArgGhost,
1598  bool const previous) { return previous; }
1599  };
1600  template<typename StencilListT, typename Dummy>
1601  struct VerifyNoGhostUseOnEdge<StencilListT,
1603  Dummy> {
1604  static inline bool verify(StencilListT const & stencils,
1605  GhostData const & baseArgGhost,
1606  bool const previous) {
1607  return VerifyNoGhostUseOnEdge<typename StencilListT::Collection,
1608  DomainEdgeSide::PLUS_SIDE>::verify(stencils.others(),
1609  baseArgGhost,
1610  previous
1611  &&
1612  baseArgGhost.get_plus(DirT::
1613  value)
1614  -
1615  stencils.current().ghosts_with_bc().get_plus(DirT::
1616  value)
1617  <
1618  StencilListT::
1619  length);
1620  }
1621  };
1622  template<typename Dummy>
1623  struct VerifyNoGhostUseOnEdge<NeboGenericEmptyTypeList,
1625  Dummy> {
1626  static inline bool verify(NeboGenericEmptyTypeList const & stencils,
1627  GhostData const & baseArgGhost,
1628  bool const previous) { return previous; }
1629  };
1630 
1631  inline GhostData varying_ghosts_with_bc(void) const {
1632  GhostData const mainStencilGhost(mainST_.ghosts_without_bc());
1633 
1634  GhostData const mainBCCells(mainST_.ghosts_with_bc() -
1635  mainStencilGhost);
1636 
1637  if(executionSide_ == static_cast<int>(DomainEdgeSide::BOTH_SIDE::
1638  value)) {
1639  /* Both sides means we do not fill in directional ghost data nor need it
1640  */
1641  IntVec plus = mainStencilGhost.get_plus();
1642  plus[DirT::value] = 0;
1643  IntVec minus = mainStencilGhost.get_minus();
1644  minus[DirT::value] = 0;
1645  return GhostData(minus, plus) + mainBCCells;
1646  }
1647  else {
1648  if(executionSide_ == static_cast<int>(DomainEdgeSide::MINUS_SIDE::
1649  value)) {
1650  IntVec minus = mainStencilGhost.get_minus();
1651 
1652  minus[DirT::value] = 0;
1653 
1654  return GhostData(minus, mainStencilGhost.get_plus()) +
1655  mainBCCells;
1656  }
1657  else if(executionSide_ == static_cast<int>(DomainEdgeSide::
1658  PLUS_SIDE::value)) {
1659  IntVec plus = mainStencilGhost.get_plus();
1660 
1661  plus[DirT::value] = 0;
1662 
1663  return GhostData(mainStencilGhost.get_minus(), plus) +
1664  mainBCCells;
1665  };
1666 
1667  return mainStencilGhost + mainBCCells;
1668  };
1669  }
1670 
1671  inline GhostData varying_ghosts_without_bc(void) const {
1672  GhostData const mainStencilGhost(mainST_.ghosts_without_bc());
1673 
1674  if(executionSide_ == static_cast<int>(DomainEdgeSide::BOTH_SIDE::
1675  value)) {
1676  /* Both sides means we do not fill in directional ghost data nor need it
1677  */
1678  IntVec plus = mainStencilGhost.get_plus();
1679  plus[DirT::value] = 0;
1680  IntVec minus = mainStencilGhost.get_minus();
1681  minus[DirT::value] = 0;
1682  return GhostData(minus, plus);
1683  }
1684  else {
1685  if(executionSide_ == static_cast<int>(DomainEdgeSide::MINUS_SIDE::
1686  value)) {
1687  IntVec minus = mainStencilGhost.get_minus();
1688 
1689  minus[DirT::value] = 0;
1690 
1691  return GhostData(minus, mainStencilGhost.get_plus());
1692  }
1693  else if(executionSide_ == static_cast<int>(DomainEdgeSide::
1694  PLUS_SIDE::value)) {
1695  IntVec plus = mainStencilGhost.get_plus();
1696 
1697  plus[DirT::value] = 0;
1698 
1699  return GhostData(mainStencilGhost.get_minus(), plus);
1700  };
1701 
1702  return mainStencilGhost;
1703  };
1704  }
1705 
1707  typename StencilT::SeqWalkType,
1708  typename ConvertToSeqWalk<NegativeStencilTs>::
1709  Result,
1710  typename ConvertToSeqWalk<PositiveStencilTs>::
1711  Result,
1712  typename Arg::SeqWalkType,
1713  DirT,
1714  FieldType> typedef SeqWalkType;
1715 
1716  #ifdef ENABLE_THREADS
1718  typename StencilT::ResizeType,
1719  typename ConvertToResize<NegativeStencilTs>::
1720  Result,
1721  typename ConvertToResize<PositiveStencilTs>::
1722  Result,
1723  typename Arg::ResizeType,
1724  DirT,
1725  FieldType> typedef ResizeType;
1726  #endif
1727  /* ENABLE_THREADS */
1728 
1729  #ifdef __CUDACC__
1731  typename StencilT::GPUWalkType,
1732  typename ConvertToGPUWalk<NegativeStencilTs>::
1733  Result,
1734  typename ConvertToGPUWalk<PositiveStencilTs>::
1735  Result,
1736  typename Arg::GPUWalkType,
1737  DirT,
1738  FieldType> typedef GPUWalkType;
1739  #endif
1740  /* __CUDACC__ */
1741 
1742  NeboVaryingEdgeStencilOneDim(Arg const & a,
1743  StencilT const & mainStencil,
1744  NegativeStencilTs const & negStencils,
1745  PositiveStencilTs const & posStencils)
1746  : arg_(a), mainST_(mainStencil), negST_(negStencils), posST_(posStencils)
1747  {
1748  /* Compute which side we are executing on */
1749  {
1750  bool const minusBC = arg_.has_bc(MINUS_SIDE)[DirT::value];
1751 
1752  bool const plusBC = arg_.has_bc(PLUS_SIDE)[DirT::value];
1753 
1754  if(minusBC && plusBC) {
1755  executionSide_ = DomainEdgeSide::BOTH_SIDE::value;
1756  }
1757  else if(minusBC) {
1758  executionSide_ = DomainEdgeSide::MINUS_SIDE::value;
1759  }
1760  else if(plusBC) {
1761  executionSide_ = DomainEdgeSide::PLUS_SIDE::value;
1762  }
1763  else { executionSide_ = DomainEdgeSide::NO_SIDE::value; };
1764  };
1765 
1766  #ifndef NDEBUG
1767  GhostData const baseArgGhost(arg_.ghosts_without_bc());
1768  GhostData const neededGhost(mainST_.ghosts_without_bc() -
1769  baseArgGhost);
1770  /* Checking enough negative stencils given based on stencil width
1771  */
1772  assert(executionSide_ == static_cast<int>(DomainEdgeSide::
1773  PLUS_SIDE::value) ||
1774  executionSide_ == static_cast<int>(DomainEdgeSide::
1775  NO_SIDE::value) ||
1776  neededGhost.get_minus(DirT::value) == -(NegativeStencilTs::
1777  length));
1778  /* Checking enough positive stencils given based on stencil width
1779  */
1780  assert(executionSide_ == static_cast<int>(DomainEdgeSide::
1781  MINUS_SIDE::value) ||
1782  executionSide_ == static_cast<int>(DomainEdgeSide::
1783  NO_SIDE::value) ||
1784  neededGhost.get_plus(DirT::value) == -(PositiveStencilTs::
1785  length));
1786  /* Checking no ghost cells used on minus side if it is a domain edge
1787  */
1788  if(executionSide_ != static_cast<int>(DomainEdgeSide::PLUS_SIDE::
1789  value) && executionSide_
1790  != static_cast<int>(DomainEdgeSide::NO_SIDE::value)) {
1791  assert((VerifyNoGhostUseOnEdge<NegativeStencilTs,
1793  verify(negST_, baseArgGhost, true)));
1794  }
1795  /* Checking no ghost cells used on plus side if it is a domain edge
1796  */
1797  if(executionSide_ != static_cast<int>(DomainEdgeSide::MINUS_SIDE::
1798  value) && executionSide_
1799  != static_cast<int>(DomainEdgeSide::NO_SIDE::value)) {
1800  assert((VerifyNoGhostUseOnEdge<PositiveStencilTs,
1802  verify(posST_, baseArgGhost, true)));
1803  }
1804  #endif
1805  /* NDEBUG */;
1806  }
1807 
1808  inline GhostData ghosts_with_bc(void) const {
1809  return varying_ghosts_with_bc();
1810  }
1811 
1812  inline GhostData ghosts_without_bc(void) const {
1813  return varying_ghosts_without_bc();
1814  }
1815 
1816  inline bool has_extents(void) const { return arg_.has_extents(); }
1817 
1818  inline IntVec extents(void) const { return arg_.extents(); }
1819 
1820  inline IntVec has_bc(BCSide const bcSide) const {
1821  return arg_.has_bc(bcSide);
1822  }
1823 
1824  inline SeqWalkType init(IntVec const & extents,
1825  GhostData const & ghosts,
1826  IntVec const & hasBC) const {
1827  return SeqWalkType(extents + ghosts.get_plus(),
1828  arg_.init(extents, ghosts, hasBC),
1829  mainST_.init(extents, ghosts, hasBC),
1830  ConvertToSeqWalk<NegativeStencilTs>::
1831  convert_runtime(negST_, extents, ghosts, hasBC),
1832  ConvertToSeqWalk<PositiveStencilTs>::
1833  convert_runtime(posST_, extents, ghosts, hasBC),
1834  executionSide_);
1835  }
1836 
1837  #ifdef ENABLE_THREADS
1838  inline ResizeType resize(void) const {
1839  return ResizeType(arg_.resize(),
1840  mainST_.resize(),
1841  ConvertToResize<NegativeStencilTs>::
1842  convert_runtime(negST_),
1843  ConvertToResize<PositiveStencilTs>::
1844  convert_runtime(posST_),
1845  executionSide_);
1846  }
1847  #endif
1848  /* ENABLE_THREADS */
1849 
1850  #ifdef __CUDACC__
1851  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
1852 
1853  inline bool gpu_ready(int const deviceIndex) const {
1854  return arg_.gpu_ready(deviceIndex);
1855  }
1856 
1857  inline GPUWalkType gpu_init(IntVec const & extents,
1858  GhostData const & ghosts,
1859  IntVec const & hasBC,
1860  int const deviceIndex,
1861  cudaStream_t const & lhsStream) const {
1862  return GPUWalkType(extents + ghosts.get_plus(),
1863  arg_.gpu_init(extents,
1864  ghosts,
1865  hasBC,
1866  deviceIndex,
1867  lhsStream),
1868  mainST_.gpu_init(extents,
1869  ghosts,
1870  hasBC,
1871  deviceIndex,
1872  lhsStream),
1873  ConvertToGPUWalk<NegativeStencilTs>::
1874  convert_runtime(negST_,
1875  extents,
1876  ghosts,
1877  hasBC,
1878  deviceIndex,
1879  lhsStream),
1880  ConvertToGPUWalk<PositiveStencilTs>::
1881  convert_runtime(posST_,
1882  extents,
1883  ghosts,
1884  hasBC,
1885  deviceIndex,
1886  lhsStream),
1887  executionSide_);
1888  }
1889 
1890  inline void stream_wait_event(cudaEvent_t const & event) const {
1891  arg_.stream_wait_event(event);
1892  }
1893 
1894  #ifdef NEBO_GPU_TEST
1895  inline void gpu_prep(int const deviceIndex) const {
1896  arg_.gpu_prep(deviceIndex);
1897  }
1898  #endif
1899  /* NEBO_GPU_TEST */
1900  #endif
1901  /* __CUDACC__ */
1902 
1903  private:
1904  Arg const arg_;
1905 
1906  StencilT const mainST_;
1907 
1908  NegativeStencilTs const negST_;
1909 
1910  PositiveStencilTs const posST_;
1911 
1912  int executionSide_;
1913  };
1914  #ifdef ENABLE_THREADS
1915  template<typename StencilT,
1916  typename NegativeStencilTs,
1917  typename PositiveStencilTs,
1918  typename Arg,
1919  typename DirT,
1920  typename FieldType>
1921  struct NeboVaryingEdgeStencilOneDim<Resize,
1922  StencilT,
1923  NegativeStencilTs,
1924  PositiveStencilTs,
1925  Arg,
1926  DirT,
1927  FieldType> {
1928  public:
1929  FieldType typedef field_type;
1930 
1931  template<typename ListT, typename Dummy = void>
1932  struct ConvertToSeqWalk {
1933  public:
1934  NeboGenericTypeList<typename ListT::First::SeqWalkType,
1935  typename ConvertToSeqWalk<typename ListT::
1936  AllButFirst>::
1937  Result> typedef Result;
1938 
1939  static inline Result convert_runtime(ListT const list,
1940  IntVec const & extents,
1941  GhostData const & ghosts,
1942  IntVec const & hasBC) {
1943  return ConvertToSeqWalk<typename ListT::AllButFirst>::
1944  convert_runtime(list.others(), extents, ghosts, hasBC)(list.current().init(extents,
1945  ghosts,
1946  hasBC));
1947  }
1948  };
1949  template<typename Dummy>
1950  struct ConvertToSeqWalk<NeboGenericEmptyTypeList, Dummy> {
1951  public:
1952  NeboGenericEmptyTypeList typedef Result;
1953 
1954  static inline Result convert_runtime(NeboGenericEmptyTypeList
1955  const list,
1956  IntVec const & extents,
1957  GhostData const & ghosts,
1958  IntVec const & hasBC) {
1959  return list;
1960  }
1961  };
1962 
1964  typename StencilT::SeqWalkType,
1965  typename ConvertToSeqWalk<NegativeStencilTs>::
1966  Result,
1967  typename ConvertToSeqWalk<PositiveStencilTs>::
1968  Result,
1969  typename Arg::SeqWalkType,
1970  DirT,
1971  FieldType> typedef SeqWalkType;
1972 
1973  NeboVaryingEdgeStencilOneDim(Arg const & a,
1974  StencilT const & mainStencil,
1975  NegativeStencilTs const & negStencils,
1976  PositiveStencilTs const & posStencils,
1977  int const executionSide)
1978  : arg_(a),
1979  mainST_(mainStencil),
1980  negST_(negStencils),
1981  posST_(posStencils),
1982  executionSide_(executionSide)
1983  {}
1984 
1985  inline SeqWalkType init(IntVec const & extents,
1986  GhostData const & ghosts,
1987  IntVec const & hasBC) const {
1988  return SeqWalkType(extents + ghosts.get_plus(),
1989  arg_.init(extents, ghosts, hasBC),
1990  mainST_.init(extents, ghosts, hasBC),
1991  ConvertToSeqWalk<NegativeStencilTs>::
1992  convert_runtime(negST_,
1993  extents,
1994  ghosts,
1995  hasBC),
1996  ConvertToSeqWalk<PositiveStencilTs>::
1997  convert_runtime(posST_,
1998  extents,
1999  ghosts,
2000  hasBC),
2001  executionSide_);
2002  }
2003 
2004  private:
2005  Arg const arg_;
2006 
2007  StencilT const mainST_;
2008 
2009  NegativeStencilTs const negST_;
2010 
2011  PositiveStencilTs const posST_;
2012 
2013  int const executionSide_;
2014  }
2015  #endif
2016  /* ENABLE_THREADS */;
2017  template<typename StencilT,
2018  typename NegativeStencilTs,
2019  typename PositiveStencilTs,
2020  typename Arg,
2021  typename DirT,
2022  typename FieldType>
2024  StencilT,
2025  NegativeStencilTs,
2026  PositiveStencilTs,
2027  Arg,
2028  DirT,
2029  FieldType> {
2030  public:
2031  FieldType typedef field_type;
2032 
2033  typename field_type::value_type typedef value_type;
2034 
2036  StencilT,
2037  NegativeStencilTs,
2038  PositiveStencilTs,
2039  Arg,
2040  DirT,
2041  FieldType> typedef MyType;
2042 
2043  template<typename ListT, int CurrentIndex>
2044  struct EvalEdgeStencil {
2045  template<typename OptionalArgT>
2046  static inline value_type eval(ListT const & list,
2047  int const searchIndex,
2048  int const x,
2049  int const y,
2050  int const z) {
2051  return (CurrentIndex == searchIndex ? list.current().template
2052  eval<OptionalArgT>(x,
2053  y,
2054  z)
2055  : EvalEdgeStencil<typename ListT::Collection,
2056  CurrentIndex - 1>::template eval<OptionalArgT>(list.others(),
2057  searchIndex,
2058  x,
2059  y,
2060  z));
2061  }
2062  };
2063  template<typename ListT>
2064  struct EvalEdgeStencil<ListT, 0> {
2065  template<typename OptionalArgT>
2066  static inline value_type eval(ListT const & list,
2067  int const searchIndex,
2068  int const x,
2069  int const y,
2070  int const z) {
2071  return list.current().template eval<OptionalArgT>(x, y, z);
2072  }
2073  };
2074 
2075  template<typename IndexDirT, typename Dummy = void>
2076  struct GetIndex;
2077  template<typename Dummy>
2078  struct GetIndex<XDIR, Dummy> {
2079  public:
2080  static inline int index(int const x, int const y, int const z) {
2081  return x;
2082  }
2083  };
2084  template<typename Dummy>
2085  struct GetIndex<YDIR, Dummy> {
2086  public:
2087  static inline int index(int const x, int const y, int const z) {
2088  return y;
2089  }
2090  };
2091  template<typename Dummy>
2092  struct GetIndex<ZDIR, Dummy> {
2093  public:
2094  static inline int index(int const x, int const y, int const z) {
2095  return z;
2096  }
2097  };
2098 
2099  template<typename T, typename Dummy = void>
2100  struct SpecialArgHandler {
2101  constexpr static inline bool hasBC(void) { return false; }
2102  };
2103  template<typename Dummy>
2104  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::HasBCOnX,
2105  Dummy> {
2106  constexpr static inline bool hasBC(void) {
2107  return std::is_same<DirT, XDIR>::value;
2108  }
2109  };
2110  template<typename Dummy>
2111  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::HasBCOnY,
2112  Dummy> {
2113  constexpr static inline bool hasBC(void) {
2114  return std::is_same<DirT, YDIR>::value;
2115  }
2116  };
2117  template<typename Dummy>
2118  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::HasBCOnZ,
2119  Dummy> {
2120  constexpr static inline bool hasBC(void) {
2121  return std::is_same<DirT, ZDIR>::value;
2122  }
2123  };
2124  template<typename ... OptionalArgT>
2125  struct OptionalArgParser;
2126  template<typename ... T>
2127  struct OptionalArgParser<CompileTimeOptionalArgs<T ...> > {
2128  constexpr static inline bool hasBC(void) {
2129  return OptionalArgParser<T ...>::hasBC();
2130  }
2131  };
2132  template<typename U, typename ... T>
2133  struct OptionalArgParser<U, T ...> {
2134  constexpr static inline bool hasBC(void) {
2135  return SpecialArgHandler<U>::hasBC() || OptionalArgParser<T ...>::
2136  hasBC();
2137  }
2138  };
2139  template<typename T>
2140  struct OptionalArgParser<T> {
2141  constexpr static inline bool hasBC(void) {
2142  return SpecialArgHandler<T>::hasBC() || false;
2143  }
2144  };
2145 
2146  template<typename OptionalArgT>
2147  inline value_type evalOnSide(int const x, int const y, int const z) const {
2148  if(OptionalArgParser<OptionalArgT>::hasBC()) {
2149  if(executionSide_ != static_cast<int>(DomainEdgeSide::NO_SIDE::
2150  value)) {
2151  int const i = GetIndex<DirT>::index(x, y, z);
2152 
2153  int const negi = (dirExtent_ - 1) - i;
2154 
2155  if(executionSide_ != static_cast<int>(DomainEdgeSide::
2156  MINUS_SIDE::value) &&
2157  0 <= negi && negi < PositiveStencilTs::length) {
2158  return EvalEdgeStencil<PositiveStencilTs,
2159  PositiveStencilTs::length - 1>::
2160  template eval<OptionalArgT>(posST_, negi, x, y, z);
2161  }
2162  else if(executionSide_ != static_cast<int>(DomainEdgeSide::
2163  PLUS_SIDE::value)
2164  && 0 <= i && i < NegativeStencilTs::length) {
2165  return EvalEdgeStencil<NegativeStencilTs,
2166  NegativeStencilTs::length - 1>::
2167  template eval<OptionalArgT>(negST_, i, x, y, z);
2168  };
2169  };
2170  };
2171 
2172  return mainST_.template eval<OptionalArgT>(x, y, z);
2173  }
2174 
2175  NeboVaryingEdgeStencilOneDim(IntVec const & plusLimit,
2176  Arg const & arg,
2177  StencilT const & mainStencil,
2178  NegativeStencilTs const & negStencils,
2179  PositiveStencilTs const & posStencils,
2180  int const executionSide)
2181  : dirExtent_(GetIndex<DirT>::index(plusLimit[0],
2182  plusLimit[1],
2183  plusLimit[2])),
2184  mainST_(mainStencil),
2185  negST_(negStencils),
2186  posST_(posStencils),
2187  executionSide_(executionSide)
2188  {}
2189 
2190  template<typename OptionalArgT>
2191  inline value_type eval(int const x, int const y, int const z) const {
2192  return evalOnSide<OptionalArgT>(x, y, z);
2193  }
2194 
2195  private:
2196  int const dirExtent_;
2197 
2198  StencilT const mainST_;
2199 
2200  NegativeStencilTs const negST_;
2201 
2202  PositiveStencilTs const posST_;
2203 
2204  int const executionSide_;
2205  };
2206  #ifdef __CUDACC__
2207  template<typename StencilT,
2208  typename NegativeStencilTs,
2209  typename PositiveStencilTs,
2210  typename Arg,
2211  typename DirT,
2212  typename FieldType>
2213  struct NeboVaryingEdgeStencilOneDim<GPUWalk,
2214  StencilT,
2215  NegativeStencilTs,
2216  PositiveStencilTs,
2217  Arg,
2218  DirT,
2219  FieldType> {
2220  public:
2221  FieldType typedef field_type;
2222 
2223  typename field_type::value_type typedef value_type;
2224 
2225  template<typename ListT, int CurrentIndex>
2226  struct EvalEdgeStencil {
2227  template<typename OptionalArgT>
2228  __device__ static inline value_type eval(ListT const & list,
2229  int const searchIndex,
2230  int const x,
2231  int const y,
2232  int const z) {
2233  return (CurrentIndex == searchIndex ? list.current().template
2234  eval<OptionalArgT>(x,
2235  y,
2236  z)
2237  : EvalEdgeStencil<typename ListT::Collection,
2238  CurrentIndex - 1>::template eval<OptionalArgT>(list.others(),
2239  searchIndex,
2240  x,
2241  y,
2242  z));
2243  }
2244  };
2245  template<typename ListT>
2246  struct EvalEdgeStencil<ListT, 0> {
2247  template<typename OptionalArgT>
2248  __device__ static inline value_type eval(ListT const & list,
2249  int const searchIndex,
2250  int const x,
2251  int const y,
2252  int const z) {
2253  return list.current().template eval<OptionalArgT>(x, y, z);
2254  }
2255  };
2256 
2257  template<typename IndexDirT, typename Dummy = void>
2258  struct GetIndex;
2259  template<typename Dummy>
2260  struct GetIndex<XDIR, Dummy> {
2261  public:
2262  #ifdef __CUDACC__
2263  __host__ __device__
2264  #endif
2265  /* __CUDACC__ */ static inline int index(int const x,
2266  int const y,
2267  int const z) {
2268  return x;
2269  }
2270  };
2271  template<typename Dummy>
2272  struct GetIndex<YDIR, Dummy> {
2273  public:
2274  #ifdef __CUDACC__
2275  __host__ __device__
2276  #endif
2277  /* __CUDACC__ */ static inline int index(int const x,
2278  int const y,
2279  int const z) {
2280  return y;
2281  }
2282  };
2283  template<typename Dummy>
2284  struct GetIndex<ZDIR, Dummy> {
2285  public:
2286  #ifdef __CUDACC__
2287  __host__ __device__
2288  #endif
2289  /* __CUDACC__ */ static inline int index(int const x,
2290  int const y,
2291  int const z) {
2292  return z;
2293  }
2294  };
2295 
2296  template<typename T, typename Dummy = void>
2297  struct SpecialArgHandler {
2298  __device__ constexpr static inline bool hasBC(void) {
2299  return false;
2300  }
2301  };
2302  template<typename Dummy>
2303  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::
2304  HasBCOnX,
2305  Dummy> {
2306  __device__ constexpr static inline bool hasBC(void) {
2307  return std::is_same<DirT, XDIR>::value;
2308  }
2309  };
2310  template<typename Dummy>
2311  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::
2312  HasBCOnY,
2313  Dummy> {
2314  __device__ constexpr static inline bool hasBC(void) {
2315  return std::is_same<DirT, YDIR>::value;
2316  }
2317  };
2318  template<typename Dummy>
2319  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::
2320  HasBCOnZ,
2321  Dummy> {
2322  __device__ constexpr static inline bool hasBC(void) {
2323  return std::is_same<DirT, ZDIR>::value;
2324  }
2325  };
2326  template<typename ... OptionalArgT>
2327  struct OptionalArgParser;
2328  template<typename ... T>
2329  struct OptionalArgParser<CompileTimeOptionalArgs<T ...> > {
2330  __device__ constexpr static inline bool hasBC(void) {
2331  return OptionalArgParser<T ...>::hasBC();
2332  }
2333  };
2334  template<typename U, typename ... T>
2335  struct OptionalArgParser<U, T ...> {
2336  __device__ constexpr static inline bool hasBC(void) {
2337  return SpecialArgHandler<U>::hasBC() || OptionalArgParser<T
2338  ...>::
2339  hasBC();
2340  }
2341  };
2342  template<typename T>
2343  struct OptionalArgParser<T> {
2344  __device__ constexpr static inline bool hasBC(void) {
2345  return SpecialArgHandler<T>::hasBC() || false;
2346  }
2347  };
2348 
2349  template<typename OptionalArgT>
2350  __device__ inline value_type evalOnSide(int const x,
2351  int const y,
2352  int const z) const {
2353  if(OptionalArgParser<OptionalArgT>::hasBC()) {
2354  if(executionSide_ != static_cast<int>(DomainEdgeSide::
2355  NO_SIDE::value)) {
2356  int const i = GetIndex<DirT>::index(x, y, z);
2357 
2358  int const negi = (dirExtent_ - 1) - i;
2359 
2360  if(executionSide_ != static_cast<int>(DomainEdgeSide::
2361  MINUS_SIDE::value)
2362  && 0 <= negi && negi < PositiveStencilTs::length) {
2363  return EvalEdgeStencil<PositiveStencilTs,
2364  PositiveStencilTs::length - 1>::
2365  template eval<OptionalArgT>(posST_,
2366  negi,
2367  x,
2368  y,
2369  z);
2370  }
2371  else if(executionSide_ != static_cast<int>(DomainEdgeSide::
2372  PLUS_SIDE::
2373  value) && 0 <=
2374  i && i < NegativeStencilTs::length) {
2375  return EvalEdgeStencil<NegativeStencilTs,
2376  NegativeStencilTs::length - 1>::
2377  template eval<OptionalArgT>(negST_, i, x, y, z);
2378  };
2379  };
2380  };
2381 
2382  return mainST_.template eval<OptionalArgT>(x, y, z);
2383  }
2384 
2385  NeboVaryingEdgeStencilOneDim(IntVec const & extents,
2386  Arg const & arg,
2387  StencilT const & mainStencil,
2388  NegativeStencilTs const & negStencils,
2389  PositiveStencilTs const & posStencils,
2390  int const executionSide)
2391  : dirExtent_(GetIndex<DirT>::index(extents[0],
2392  extents[1],
2393  extents[2])),
2394  mainST_(mainStencil),
2395  negST_(negStencils),
2396  posST_(posStencils),
2397  executionSide_(executionSide)
2398  {}
2399 
2400  template<typename OptionalArgT>
2401  __device__ inline value_type eval(int const x,
2402  int const y,
2403  int const z) const {
2404  return evalOnSide<OptionalArgT>(x, y, z);
2405  }
2406 
2407  private:
2408  int const dirExtent_;
2409 
2410  StencilT const mainST_;
2411 
2412  NegativeStencilTs const negST_;
2413 
2414  PositiveStencilTs const posST_;
2415 
2416  int const executionSide_;
2417  }
2418  #endif
2419  /* __CUDACC__ */;
2420  } /* SpatialOps */
2421 
2422 #endif
2423 /* NEBO_STENCILS_H */
Defines a domain edge on the minus side of a dimension.
Defines a domain edge on the plus side of a dimension.
Defines no domain edge on a dimension.
Defines a domain edge on both sides of a dimension.
IntVec get_plus() const
obtain the IntVec containing the number of ghost cells on the (+) faces
Definition: GhostData.h:145
Holds information about the number of ghost cells on each side of the domain.
Definition: GhostData.h:54
BCSide
Allows identification of whether we are setting the BC on the right or left side when using an operat...
Defines a type for the x-direction.
Defines a type for the y-direction.
Defines a type for the z-direction.
IntVec get_minus() const
obtain the IntVec containing the number of ghost cells on the (-) faces
Definition: GhostData.h:135
Allows identification of whether a domain edge lies on the minus, positive, or both sides of a single...