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,
482  NeboOptionalArg & optArg) const {
483  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg), coefs_);
484  }
485 
486  #ifdef ENABLE_THREADS
487  inline ResizeType resize(void) const {
488  return ResizeType(arg_.resize(), coefs_);
489  }
490  #endif
491  /* ENABLE_THREADS */
492 
493  #ifdef __CUDACC__
494  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
495 
496  inline bool gpu_ready(int const deviceIndex) const {
497  return arg_.gpu_ready(deviceIndex);
498  }
499 
500  inline GPUWalkType gpu_init(IntVec const & extents,
501  GhostData const & ghosts,
502  IntVec const & hasBC,
503  int const deviceIndex,
504  cudaStream_t const & lhsStream,
505  NeboOptionalArg & optArg) const {
506  return GPUWalkType(arg_.gpu_init(extents,
507  ghosts,
508  hasBC,
509  deviceIndex,
510  lhsStream,
511  optArg),
512  coefs_);
513  }
514 
515  inline void stream_wait_event(cudaEvent_t const & event) const {
516  arg_.stream_wait_event(event);
517  }
518 
519  #ifdef NEBO_GPU_TEST
520  inline void gpu_prep(int const deviceIndex) const {
521  arg_.gpu_prep(deviceIndex);
522  }
523  #endif
524  /* NEBO_GPU_TEST */
525  #endif
526  /* __CUDACC__ */
527 
528  private:
529  Arg const arg_;
530 
531  Coefs const coefs_;
532  };
533  #ifdef ENABLE_THREADS
534  template<typename Pts, typename Arg, typename FieldType>
535  struct NeboStencil<Resize, Pts, Arg, FieldType> {
536  public:
537  FieldType typedef field_type;
538 
540 
542  typedef SeqWalkType;
543 
544  NeboStencil(Arg const & arg, Coefs const & coefs)
545  : arg_(arg), coefs_(coefs)
546  {}
547 
548  inline SeqWalkType init(IntVec const & extents,
549  GhostData const & ghosts,
550  IntVec const & hasBC,
551  NeboOptionalArg & optArg) const {
552  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg), coefs_);
553  }
554 
555  private:
556  Arg const arg_;
557 
558  Coefs const coefs_;
559  }
560  #endif
561  /* ENABLE_THREADS */;
562  template<typename Pts, typename Arg, typename FieldType>
563  struct NeboStencil<SeqWalk, Pts, Arg, FieldType> {
564  public:
565  FieldType typedef field_type;
566 
567  typename field_type::value_type typedef value_type;
568 
570 
571  template<typename PointCollection>
572  struct EvalExpr {
574 
575  typename PointCollection::Point typedef Point;
576 
577  typename PointCollection::Collection typedef Collection;
578 
579  template<typename OptionalArgT>
580  static inline value_type eval(Arg const & arg,
581  Coefs const & coefs,
582  int const x,
583  int const y,
584  int const z) {
585  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
586  coefs.others(),
587  x,
588  y,
589  z) +
590  arg.template eval<OptionalArgT>(x + Point::value(0),
591  y + Point::value(1),
592  z + Point::value(2)) *
593  coefs.coef();
594  }
595  };
596 
597  template<typename Point>
598  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
599  NeboStencilCoefCollection<1> typedef Coefs;
600 
601  template<typename OptionalArgT>
602  static inline value_type eval(Arg const & arg,
603  Coefs const & coefs,
604  int const x,
605  int const y,
606  int const z) {
607  return arg.template eval<OptionalArgT>(x + Point::value(0),
608  y + Point::value(1),
609  z + Point::value(2)) *
610  coefs.coef();
611  }
612  };
613 
614  NeboStencil(Arg const & arg, Coefs const & coefs)
615  : arg_(arg), coefs_(coefs)
616  {}
617 
618  template<typename OptionalArgT>
619  inline value_type eval(int const x, int const y, int const z) const {
620  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
621  coefs_,
622  x,
623  y,
624  z);
625  }
626 
627  private:
628  Arg arg_;
629 
630  Coefs const coefs_;
631  };
632  #ifdef __CUDACC__
633  template<typename Pts, typename Arg, typename FieldType>
634  struct NeboStencil<GPUWalk, Pts, Arg, FieldType> {
635  public:
636  FieldType typedef field_type;
637 
638  typename field_type::value_type typedef value_type;
639 
641 
642  template<typename PointCollection>
643  struct EvalExpr {
645  ;
646 
647  typename PointCollection::Point typedef Point;
648 
649  typename PointCollection::Collection typedef Collection;
650 
651  template<typename OptionalArgT>
652  __device__ static inline value_type eval(Arg const & arg,
653  Coefs const & coefs,
654  int const x,
655  int const y,
656  int const z) {
657  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
658  coefs.others(),
659  x,
660  y,
661  z)
662  + arg.template eval<OptionalArgT>(x + Point::
663  value_gpu(0),
664  y + Point::
665  value_gpu(1),
666  z + Point::
667  value_gpu(2)) *
668  coefs.coef();
669  }
670  };
671 
672  template<typename Point>
673  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
674  NeboStencilCoefCollection<1> typedef Coefs;
675 
676  template<typename OptionalArgT>
677  __device__ static inline value_type eval(Arg const & arg,
678  Coefs const & coefs,
679  int const x,
680  int const y,
681  int const z) {
682  return arg.template eval<OptionalArgT>(x + Point::value_gpu(0),
683  y + Point::value_gpu(1),
684  z + Point::value_gpu(2))
685  * coefs.coef();
686  }
687  };
688 
689  NeboStencil(Arg const & a, Coefs const & coefs)
690  : arg_(a), coefs_(coefs)
691  {}
692 
693  template<typename OptionalArgT>
694  __device__ inline value_type eval(int const x,
695  int const y,
696  int const z) const {
697  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
698  coefs_,
699  x,
700  y,
701  z);
702  }
703 
704  private:
705  Arg arg_;
706 
707  Coefs const coefs_;
708  }
709  #endif
710  /* __CUDACC__ */;
711 
712  template<typename CurrentMode,
713  typename Pts,
714  typename Arg,
715  typename FieldType>
717  template<typename Pts, typename Arg, typename FieldType>
718  struct NeboEdgelessStencil<Initial, Pts, Arg, FieldType> {
719  public:
720  FieldType typedef field_type;
721 
723 
725  typedef SeqWalkType;
726 
727  #ifdef ENABLE_THREADS
728  NeboEdgelessStencil<Resize,
729  Pts,
730  typename Arg::ResizeType,
731  FieldType> typedef ResizeType;
732  #endif
733  /* ENABLE_THREADS */
734 
735  #ifdef __CUDACC__
736  NeboEdgelessStencil<GPUWalk,
737  Pts,
738  typename Arg::GPUWalkType,
739  FieldType> typedef GPUWalkType;
740  #endif
741  /* __CUDACC__ */
742 
743  NeboEdgelessStencil(Arg const & a, Coefs const & coefs)
744  : arg_(a), coefs_(coefs)
745  {}
746 
747  inline GhostData ghosts_with_bc(void) const {
748  return Pts::possible_additive_ghosts(arg_.ghosts_with_bc());
749  }
750 
751  inline GhostData ghosts_without_bc(void) const {
752  return Pts::possible_additive_ghosts(arg_.ghosts_without_bc());
753  }
754 
755  inline bool has_extents(void) const { return arg_.has_extent(); }
756 
757  inline IntVec extents(void) const { return arg_.extents(); }
758 
759  inline IntVec has_bc(BCSide const bcSide) const {
760  return arg_.has_bc(bcSide);
761  }
762 
763  inline SeqWalkType init(IntVec const & extents,
764  GhostData const & ghosts,
765  IntVec const & hasBC,
766  NeboOptionalArg & optArg) const {
767  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg),
768  coefs_,
769  lowest_indicies(),
770  highest_indicies());
771  }
772 
773  #ifdef ENABLE_THREADS
774  inline ResizeType resize(void) const {
775  return ResizeType(arg_.resize(),
776  coefs_,
777  lowest_indicies(),
778  highest_indicies());
779  }
780  #endif
781  /* ENABLE_THREADS */
782 
783  #ifdef __CUDACC__
784  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
785 
786  inline bool gpu_ready(int const deviceIndex) const {
787  return arg_.gpu_ready(deviceIndex);
788  }
789 
790  inline GPUWalkType gpu_init(IntVec const & extents,
791  GhostData const & ghosts,
792  IntVec const & hasBC,
793  int const deviceIndex,
794  cudaStream_t const & lhsStream,
795  NeboOptionalArg & optArg) const {
796  return GPUWalkType(arg_.gpu_init(extents,
797  ghosts,
798  hasBC,
799  deviceIndex,
800  lhsStream,
801  optArg),
802  coefs_);
803  }
804 
805  inline void stream_wait_event(cudaEvent_t const & event) const {
806  arg_.stream_wait_event(event);
807  }
808 
809  #ifdef NEBO_GPU_TEST
810  inline void gpu_prep(int const deviceIndex) const {
811  arg_.gpu_prep(deviceIndex);
812  }
813  #endif
814  /* NEBO_GPU_TEST */
815  #endif
816  /* __CUDACC__ */
817 
818  inline GhostData actual_ghosts(void) const {
819  return Pts::possible_ghosts(arg_.ghosts_with_bc());
820  }
821 
822  inline IntVec lowest_indicies(void) const {
823  return -(actual_ghosts().get_minus());
824  }
825 
826  inline IntVec highest_indicies(void) const {
827  return actual_ghosts().get_plus() + extents();
828  }
829 
830  private:
831  Arg const arg_;
832 
833  Coefs const coefs_;
834  };
835  #ifdef ENABLE_THREADS
836  template<typename Pts, typename Arg, typename FieldType>
837  struct NeboEdgelessStencil<Resize, Pts, Arg, FieldType> {
838  public:
839  FieldType typedef field_type;
840 
842 
843  NeboEdgelessStencil<SeqWalk,
844  Pts,
845  typename Arg::SeqWalkType,
846  FieldType> typedef SeqWalkType;
847 
848  NeboEdgelessStencil(Arg const & arg,
849  Coefs const & coefs,
850  IntVec const & low,
851  IntVec const & high)
852  : arg_(arg), coefs_(coefs), low_(low), high_(high)
853  {}
854 
855  inline SeqWalkType init(IntVec const & extents,
856  GhostData const & ghosts,
857  IntVec const & hasBC,
858  NeboOptionalArg & optArg) const {
859  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg),
860  coefs_,
861  low_,
862  high_);
863  }
864 
865  private:
866  Arg const arg_;
867 
868  Coefs const coefs_;
869 
870  IntVec const low_;
871 
872  IntVec const high_;
873  }
874  #endif
875  /* ENABLE_THREADS */;
876  template<typename Pts, typename Arg, typename FieldType>
877  struct NeboEdgelessStencil<SeqWalk, Pts, Arg, FieldType> {
878  public:
879  FieldType typedef field_type;
880 
881  typename field_type::value_type typedef value_type;
882 
884 
885  template<typename PointCollection>
886  struct EvalExpr {
888 
889  typename PointCollection::Point typedef Point;
890 
891  typename PointCollection::Collection typedef Collection;
892 
893  template<typename OptionalArgT>
894  static inline value_type eval(Arg const & arg,
895  Coefs const & coefs,
896  int const x,
897  int const y,
898  int const z) {
899  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
900  coefs.others(),
901  x,
902  y,
903  z) +
904  arg.template eval<OptionalArgT>(x + Point::value(0),
905  y + Point::value(1),
906  z + Point::value(2)) *
907  coefs.coef();
908  }
909  };
910 
911  template<typename Point>
912  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
913  NeboStencilCoefCollection<1> typedef Coefs;
914 
915  template<typename OptionalArgT>
916  static inline value_type eval(Arg const & arg,
917  Coefs const & coefs,
918  int const x,
919  int const y,
920  int const z) {
921  return arg.template eval<OptionalArgT>(x + Point::value(0),
922  y + Point::value(1),
923  z + Point::value(2)) *
924  coefs.coef();
925  }
926  };
927 
928  NeboEdgelessStencil(Arg const & arg,
929  Coefs const & coefs,
930  IntVec const & low,
931  IntVec const & high)
932  : arg_(arg), coefs_(coefs), low_(low), high_(high)
933  {}
934 
935  template<typename OptionalArgT>
936  inline value_type eval(int const x, int const y, int const z) const {
937  #ifndef NDEBUG
938  IntVec index = IntVec(x, y, z);
939  if(index < low_ || index >= high_) {
940  std::ostringstream msg;
941  msg << "Nebo error in " << "Nebo Edgeless Stencil" << ":\n";
942  msg << " - " << low_ << " < " << index << " <= " << high_;
943  msg << "\n";
944  msg << "\t - " << __FILE__ << " : " << __LINE__;
945  throw(std::runtime_error(msg.str()));;
946  }
947  #endif
948  /* NDEBUG */;
949 
950  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
951  coefs_,
952  x,
953  y,
954  z);
955  }
956 
957  private:
958  Arg arg_;
959 
960  Coefs const coefs_;
961 
962  IntVec const low_;
963 
964  IntVec const high_;
965  };
966  #ifdef __CUDACC__
967  template<typename Pts, typename Arg, typename FieldType>
968  struct NeboEdgelessStencil<GPUWalk, Pts, Arg, FieldType> {
969  public:
970  FieldType typedef field_type;
971 
972  typename field_type::value_type typedef value_type;
973 
975 
976  template<typename PointCollection>
977  struct EvalExpr {
979  ;
980 
981  typename PointCollection::Point typedef Point;
982 
983  typename PointCollection::Collection typedef Collection;
984 
985  template<typename OptionalArgT>
986  __device__ static inline value_type eval(Arg const & arg,
987  Coefs const & coefs,
988  int const x,
989  int const y,
990  int const z) {
991  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
992  coefs.others(),
993  x,
994  y,
995  z)
996  + arg.template eval<OptionalArgT>(x + Point::
997  value_gpu(0),
998  y + Point::
999  value_gpu(1),
1000  z + Point::
1001  value_gpu(2)) *
1002  coefs.coef();
1003  }
1004  };
1005 
1006  template<typename Point>
1007  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
1008  NeboStencilCoefCollection<1> typedef Coefs;
1009 
1010  template<typename OptionalArgT>
1011  __device__ static inline value_type eval(Arg const & arg,
1012  Coefs const & coefs,
1013  int const x,
1014  int const y,
1015  int const z) {
1016  return arg.template eval<OptionalArgT>(x + Point::value_gpu(0),
1017  y + Point::value_gpu(1),
1018  z + Point::value_gpu(2))
1019  * coefs.coef();
1020  }
1021  };
1022 
1023  NeboEdgelessStencil(Arg const & a, Coefs const & coefs)
1024  : arg_(a), coefs_(coefs)
1025  {}
1026 
1027  template<typename OptionalArgT>
1028  __device__ inline value_type eval(int const x,
1029  int const y,
1030  int const z) const {
1031  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_,
1032  coefs_,
1033  x,
1034  y,
1035  z);
1036  }
1037 
1038  private:
1039  Arg arg_;
1040 
1041  Coefs const coefs_;
1042  }
1043  #endif
1044  /* __CUDACC__ */;
1045 
1046  template<typename CurrentMode,
1047  typename Pts,
1048  typename Arg,
1049  typename FieldType>
1051  template<typename Pts, typename Arg, typename FieldType>
1052  struct NeboSumStencil<Initial, Pts, Arg, FieldType> {
1053  public:
1054  FieldType typedef field_type;
1055 
1057  typedef SeqWalkType;
1058 
1059  #ifdef ENABLE_THREADS
1061  typedef ResizeType;
1062  #endif
1063  /* ENABLE_THREADS */
1064 
1065  #ifdef __CUDACC__
1067  typedef GPUWalkType;
1068  #endif
1069  /* __CUDACC__ */
1070 
1071  NeboSumStencil(Arg const & a)
1072  : arg_(a)
1073  {}
1074 
1075  inline GhostData ghosts_with_bc(void) const {
1076  return Pts::possible_ghosts(arg_.ghosts_with_bc());
1077  }
1078 
1079  inline GhostData ghosts_without_bc(void) const {
1080  return Pts::possible_ghosts(arg_.ghosts_without_bc());
1081  }
1082 
1083  inline bool has_extents(void) const { return arg_.has_extents(); }
1084 
1085  inline IntVec extents(void) const { return arg_.extents(); }
1086 
1087  inline IntVec has_bc(BCSide const bcSide) const {
1088  return arg_.has_bc(bcSide);
1089  }
1090 
1091  inline SeqWalkType init(IntVec const & extents,
1092  GhostData const & ghosts,
1093  IntVec const & hasBC,
1094  NeboOptionalArg & optArg) const {
1095  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg));
1096  }
1097 
1098  #ifdef ENABLE_THREADS
1099  inline ResizeType resize(void) const {
1100  return ResizeType(arg_.resize());
1101  }
1102  #endif
1103  /* ENABLE_THREADS */
1104 
1105  #ifdef __CUDACC__
1106  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
1107 
1108  inline bool gpu_ready(int const deviceIndex) const {
1109  return arg_.gpu_ready(deviceIndex);
1110  }
1111 
1112  inline GPUWalkType gpu_init(IntVec const & extents,
1113  GhostData const & ghosts,
1114  IntVec const & hasBC,
1115  int const deviceIndex,
1116  cudaStream_t const & lhsStream,
1117  NeboOptionalArg & optArg) const {
1118  return GPUWalkType(arg_.gpu_init(extents,
1119  ghosts,
1120  hasBC,
1121  deviceIndex,
1122  lhsStream,
1123  optArg));
1124  }
1125 
1126  inline void stream_wait_event(cudaEvent_t const & event) const {
1127  arg_.stream_wait_event(event);
1128  }
1129 
1130  #ifdef NEBO_GPU_TEST
1131  inline void gpu_prep(int const deviceIndex) const {
1132  arg_.gpu_prep(deviceIndex);
1133  }
1134  #endif
1135  /* NEBO_GPU_TEST */
1136  #endif
1137  /* __CUDACC__ */
1138 
1139  private:
1140  Arg const arg_;
1141  };
1142  #ifdef ENABLE_THREADS
1143  template<typename Pts, typename Arg, typename FieldType>
1144  struct NeboSumStencil<Resize, Pts, Arg, FieldType> {
1145  public:
1146  FieldType typedef field_type;
1147 
1149  typedef SeqWalkType;
1150 
1151  NeboSumStencil(Arg const & arg)
1152  : arg_(arg)
1153  {}
1154 
1155  inline SeqWalkType init(IntVec const & extents,
1156  GhostData const & ghosts,
1157  IntVec const & hasBC,
1158  NeboOptionalArg & optArg) const {
1159  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg));
1160  }
1161 
1162  private:
1163  Arg const arg_;
1164  }
1165  #endif
1166  /* ENABLE_THREADS */;
1167  template<typename Pts, typename Arg, typename FieldType>
1168  struct NeboSumStencil<SeqWalk, Pts, Arg, FieldType> {
1169  public:
1170  FieldType typedef field_type;
1171 
1172  typename field_type::value_type typedef value_type;
1173 
1174  template<typename PointCollection>
1175  struct EvalExpr {
1176  typename PointCollection::Point typedef Point;
1177 
1178  typename PointCollection::Collection typedef Collection;
1179 
1180  template<typename OptionalArgT>
1181  static inline value_type eval(Arg const & arg,
1182  int const x,
1183  int const y,
1184  int const z) {
1185  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
1186  x,
1187  y,
1188  z) +
1189  arg.template eval<OptionalArgT>(x + Point::value(0),
1190  y + Point::value(1),
1191  z + Point::value(2));
1192  }
1193  };
1194 
1195  template<typename Point>
1196  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
1197  template<typename OptionalArgT>
1198  static inline value_type eval(Arg const & arg,
1199  int const x,
1200  int const y,
1201  int const z) {
1202  return arg.template eval<OptionalArgT>(x + Point::value(0),
1203  y + Point::value(1),
1204  z + Point::value(2));
1205  }
1206  };
1207 
1208  NeboSumStencil(Arg const & arg)
1209  : arg_(arg)
1210  {}
1211 
1212  template<typename OptionalArgT>
1213  inline value_type eval(int const x, int const y, int const z) const {
1214  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_, x, y, z);
1215  }
1216 
1217  private:
1218  Arg arg_;
1219  };
1220  #ifdef __CUDACC__
1221  template<typename Pts, typename Arg, typename FieldType>
1222  struct NeboSumStencil<GPUWalk, Pts, Arg, FieldType> {
1223  public:
1224  FieldType typedef field_type;
1225 
1226  typename field_type::value_type typedef value_type;
1227 
1228  template<typename PointCollection>
1229  struct EvalExpr {
1230  typename PointCollection::Point typedef Point;
1231 
1232  typename PointCollection::Collection typedef Collection;
1233 
1234  template<typename OptionalArgT>
1235  __device__ static inline value_type eval(Arg const & arg,
1236  int const x,
1237  int const y,
1238  int const z) {
1239  return EvalExpr<Collection>::template eval<OptionalArgT>(arg,
1240  x,
1241  y,
1242  z)
1243  + arg.template eval<OptionalArgT>(x + Point::
1244  value_gpu(0),
1245  y + Point::
1246  value_gpu(1),
1247  z + Point::
1248  value_gpu(2));
1249  }
1250  };
1251 
1252  template<typename Point>
1253  struct EvalExpr<NeboStencilPointCollection<Point, NeboNil> > {
1254  template<typename OptionalArgT>
1255  __device__ static inline value_type eval(Arg const & arg,
1256  int const x,
1257  int const y,
1258  int const z) {
1259  return arg.template eval<OptionalArgT>(x + Point::value_gpu(0),
1260  y + Point::value_gpu(1),
1261  z + Point::value_gpu(2));
1262  }
1263  };
1264 
1265  NeboSumStencil(Arg const & a)
1266  : arg_(a)
1267  {}
1268 
1269  template<typename OptionalArgT>
1270  __device__ inline value_type eval(int const x,
1271  int const y,
1272  int const z) const {
1273  return EvalExpr<Pts>::template eval<OptionalArgT>(arg_, x, y, z);
1274  }
1275 
1276  private:
1277  Arg arg_;
1278  }
1279  #endif
1280  /* __CUDACC__ */;
1281 
1282  template<typename Point>
1283  static inline GhostData point_possible_ghosts(GhostData const & ghosts) {
1284  return ghosts + additive_reductive_point_to_ghost(Point::int_vec());
1285  };
1286 
1287  template<typename CurrentMode,
1288  typename Point,
1289  typename Arg,
1290  typename FieldType>
1292  template<typename Point, typename Arg, typename FieldType>
1293  struct NeboMaskShift<Initial, Point, Arg, FieldType> {
1294  public:
1295  FieldType typedef field_type;
1296 
1297  typename Arg::SeqWalkType typedef ArgSeqWalkType;
1298 
1299  #ifdef __CUDACC__
1300  typename Arg::GPUWalkType typedef ArgGPUWalkType;
1301  #endif
1302  /* __CUDACC__ */
1303 
1305  SeqWalkType;
1306 
1307  #ifdef ENABLE_THREADS
1309  typedef ResizeType;
1310  #endif
1311  /* ENABLE_THREADS */
1312 
1313  #ifdef __CUDACC__
1315  GPUWalkType;
1316  #endif
1317  /* __CUDACC__ */
1318 
1319  NeboMaskShift(Arg const & a)
1320  : arg_(a)
1321  {}
1322 
1323  inline GhostData ghosts_with_bc(void) const {
1324  return point_possible_ghosts<Point>(arg_.ghosts_with_bc());
1325  }
1326 
1327  inline GhostData ghosts_without_bc(void) const {
1328  return point_possible_ghosts<Point>(arg_.ghosts_without_bc());
1329  }
1330 
1331  inline bool has_extents(void) const { return arg_.has_extents(); }
1332 
1333  inline IntVec extents(void) const { return arg_.extents(); }
1334 
1335  inline IntVec has_bc(BCSide const bcSide) const {
1336  return arg_.has_bc(bcSide);
1337  }
1338 
1339  inline SeqWalkType init(IntVec const & extents,
1340  GhostData const & ghosts,
1341  IntVec const & hasBC,
1342  NeboOptionalArg & optArg) const {
1343  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg));
1344  }
1345 
1346  #ifdef ENABLE_THREADS
1347  inline ResizeType resize(void) const {
1348  return ResizeType(arg_.resize());
1349  }
1350  #endif
1351  /* ENABLE_THREADS */
1352 
1353  #ifdef __CUDACC__
1354  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
1355 
1356  inline bool gpu_ready(int const deviceIndex) const {
1357  return arg_.gpu_ready(deviceIndex);
1358  }
1359 
1360  inline GPUWalkType gpu_init(IntVec const & extents,
1361  GhostData const & ghosts,
1362  IntVec const & hasBC,
1363  int const deviceIndex,
1364  cudaStream_t const & lhsStream,
1365  NeboOptionalArg & optArg) const {
1366  return GPUWalkType(arg_.gpu_init(extents,
1367  ghosts,
1368  hasBC,
1369  deviceIndex,
1370  lhsStream,
1371  optArg));
1372  }
1373 
1374  inline void stream_wait_event(cudaEvent_t const & event) const {
1375  arg_.stream_wait_event(event);
1376  }
1377 
1378  #ifdef NEBO_GPU_TEST
1379  inline void gpu_prep(int const deviceIndex) const {
1380  arg_.gpu_prep(deviceIndex);
1381  }
1382  #endif
1383  /* NEBO_GPU_TEST */
1384  #endif
1385  /* __CUDACC__ */
1386 
1387  private:
1388  Arg const arg_;
1389  };
1390  #ifdef ENABLE_THREADS
1391  template<typename Point, typename Arg, typename FieldType>
1392  struct NeboMaskShift<Resize, Point, Arg, FieldType> {
1393  public:
1394  FieldType typedef field_type;
1395 
1396  typename Arg::SeqWalkType typedef ArgSeqWalkType;
1397 
1399  SeqWalkType;
1400 
1401  NeboMaskShift(Arg const & arg)
1402  : arg_(arg)
1403  {}
1404 
1405  inline SeqWalkType init(IntVec const & extents,
1406  GhostData const & ghosts,
1407  IntVec const & hasBC,
1408  NeboOptionalArg & optArg) const {
1409  return SeqWalkType(arg_.init(extents, ghosts, hasBC, optArg));
1410  }
1411 
1412  private:
1413  Arg const arg_;
1414  }
1415  #endif
1416  /* ENABLE_THREADS */;
1417  template<typename Point, typename Arg, typename FieldType>
1418  struct NeboMaskShift<SeqWalk, Point, Arg, FieldType> {
1419  public:
1420  FieldType typedef field_type;
1421 
1422  typename field_type::value_type typedef value_type;
1423 
1424  NeboMaskShift(Arg const & arg)
1425  : arg_(arg)
1426  {}
1427 
1428  template<typename OptionalArgT>
1429  inline bool eval(int const x, int const y, int const z) const {
1430  return arg_.template eval<OptionalArgT>(x + Point::value(0),
1431  y + Point::value(1),
1432  z + Point::value(2));
1433  }
1434 
1435  private:
1436  Arg arg_;
1437  };
1438  #ifdef __CUDACC__
1439  template<typename Point, typename Arg, typename FieldType>
1440  struct NeboMaskShift<GPUWalk, Point, Arg, FieldType> {
1441  public:
1442  FieldType typedef field_type;
1443 
1444  typename field_type::value_type typedef value_type;
1445 
1446  NeboMaskShift(Arg const & a)
1447  : arg_(a)
1448  {}
1449 
1450  template<typename OptionalArgT>
1451  __device__ inline bool eval(int const x, int const y, int const z) const {
1452  return arg_.template eval<OptionalArgT>(x + Point::value_gpu(0),
1453  y + Point::value_gpu(1),
1454  z + Point::value_gpu(2));
1455  }
1456 
1457  private:
1458  Arg arg_;
1459  }
1460  #endif
1461  /* __CUDACC__ */;
1462 
1463  template<typename CurrentMode,
1464  typename StencilT,
1465  typename NegativeStencilTs,
1466  typename PositiveStencilTs,
1467  typename Arg,
1468  typename DirT,
1469  typename FieldType>
1471  template<typename StencilT,
1472  typename NegativeStencilTs,
1473  typename PositiveStencilTs,
1474  typename Arg,
1475  typename DirT,
1476  typename FieldType>
1478  StencilT,
1479  NegativeStencilTs,
1480  PositiveStencilTs,
1481  Arg,
1482  DirT,
1483  FieldType> {
1484  public:
1485  FieldType typedef field_type;
1486 
1487  template<typename ListT, typename Dummy = void>
1488  struct ConvertToSeqWalk {
1489  public:
1490  NeboGenericTypeList<typename ListT::First::SeqWalkType,
1491  typename ConvertToSeqWalk<typename ListT::
1492  AllButFirst>::Result>
1493  typedef Result;
1494 
1495  static inline Result convert_runtime(ListT const list,
1496  IntVec const & extents,
1497  GhostData const & ghosts,
1498  IntVec const & hasBC,
1499  NeboOptionalArg & optArg) {
1500  return ConvertToSeqWalk<typename ListT::AllButFirst>::
1501  convert_runtime(list.others(), extents, ghosts, hasBC, optArg)(list.current().init(extents,
1502  ghosts,
1503  hasBC,
1504  optArg));
1505  }
1506  };
1507  template<typename Dummy>
1508  struct ConvertToSeqWalk<NeboGenericEmptyTypeList, Dummy> {
1509  public:
1510  NeboGenericEmptyTypeList typedef Result;
1511 
1512  static inline Result convert_runtime(NeboGenericEmptyTypeList
1513  const list,
1514  IntVec const & extents,
1515  GhostData const & ghosts,
1516  IntVec const & hasBC,
1517  NeboOptionalArg & optArg) {
1518  return list;
1519  }
1520  };
1521 
1522  template<typename ListT, typename Dummy = void>
1523  struct ConvertToResize {
1524  public:
1525  NeboGenericTypeList<typename ListT::First::ResizeType,
1526  typename ConvertToResize<typename ListT::
1527  AllButFirst>::Result>
1528  typedef Result;
1529 
1530  static inline Result convert_runtime(ListT const list) {
1531  return ConvertToResize<typename ListT::AllButFirst>::
1532  convert_runtime(list.others())(list.current().resize());
1533  }
1534  };
1535  template<typename Dummy>
1536  struct ConvertToResize<NeboGenericEmptyTypeList, Dummy> {
1537  public:
1538  NeboGenericEmptyTypeList typedef Result;
1539 
1540  static inline Result convert_runtime(NeboGenericEmptyTypeList
1541  const list) { return list; }
1542  };
1543 
1544  #ifdef __CUDACC__
1545  template<typename ListT, typename Dummy = void>
1546  struct ConvertToGPUWalk {
1547  public:
1548  NeboGenericTypeList<typename ListT::First::GPUWalkType,
1549  typename ConvertToGPUWalk<typename ListT::
1550  AllButFirst>::
1551  Result> typedef Result;
1552 
1553  static inline Result convert_runtime(ListT const list,
1554  IntVec const & extents,
1555  GhostData const & ghosts,
1556  IntVec const & hasBC,
1557  int const deviceIndex,
1558  cudaStream_t const &
1559  lhsStream,
1560  NeboOptionalArg & optArg) {
1561  return ConvertToGPUWalk<typename ListT::AllButFirst>::
1562  convert_runtime(list.others(),
1563  extents,
1564  ghosts,
1565  hasBC,
1566  deviceIndex,
1567  lhsStream,
1568  optArg)(list.current().gpu_init(extents,
1569  ghosts,
1570  hasBC,
1571  deviceIndex,
1572  lhsStream,
1573  optArg));
1574  }
1575  };
1576  template<typename Dummy>
1577  struct ConvertToGPUWalk<NeboGenericEmptyTypeList, Dummy> {
1578  public:
1579  NeboGenericEmptyTypeList typedef Result;
1580 
1581  static inline Result convert_runtime(NeboGenericEmptyTypeList
1582  const list,
1583  IntVec const & extents,
1584  GhostData const & ghosts,
1585  IntVec const & hasBC,
1586  int const deviceIndex,
1587  cudaStream_t const &
1588  lhsStream,
1589  NeboOptionalArg & optArg) { return list; }
1590  };
1591  #endif
1592  /* __CUDACC__ */
1593 
1594  template<typename StencilListT,
1595  typename MinusOrPlusT,
1596  typename Dummy = void>
1597  struct VerifyNoGhostUseOnEdge;
1598  template<typename StencilListT, typename Dummy>
1599  struct VerifyNoGhostUseOnEdge<StencilListT,
1601  Dummy> {
1602  static inline bool verify(StencilListT const & stencils,
1603  GhostData const & baseArgGhost,
1604  bool const previous) {
1605  return VerifyNoGhostUseOnEdge<typename StencilListT::Collection,
1607  verify(stencils.others(),
1608  baseArgGhost,
1609  previous && baseArgGhost.get_minus(DirT::value) -
1610  stencils.current().ghosts_with_bc().get_minus(DirT::
1611  value)
1612  < StencilListT::length);
1613  }
1614  };
1615  template<typename Dummy>
1616  struct VerifyNoGhostUseOnEdge<NeboGenericEmptyTypeList,
1618  Dummy> {
1619  static inline bool verify(NeboGenericEmptyTypeList const & stencils,
1620  GhostData const & baseArgGhost,
1621  bool const previous) { return previous; }
1622  };
1623  template<typename StencilListT, typename Dummy>
1624  struct VerifyNoGhostUseOnEdge<StencilListT,
1626  Dummy> {
1627  static inline bool verify(StencilListT const & stencils,
1628  GhostData const & baseArgGhost,
1629  bool const previous) {
1630  return VerifyNoGhostUseOnEdge<typename StencilListT::Collection,
1631  DomainEdgeSide::PLUS_SIDE>::verify(stencils.others(),
1632  baseArgGhost,
1633  previous
1634  &&
1635  baseArgGhost.get_plus(DirT::
1636  value)
1637  -
1638  stencils.current().ghosts_with_bc().get_plus(DirT::
1639  value)
1640  <
1641  StencilListT::
1642  length);
1643  }
1644  };
1645  template<typename Dummy>
1646  struct VerifyNoGhostUseOnEdge<NeboGenericEmptyTypeList,
1648  Dummy> {
1649  static inline bool verify(NeboGenericEmptyTypeList const & stencils,
1650  GhostData const & baseArgGhost,
1651  bool const previous) { return previous; }
1652  };
1653 
1654  inline GhostData varying_ghosts_with_bc(void) const {
1655  GhostData const mainStencilGhost(mainST_.ghosts_without_bc());
1656 
1657  GhostData const mainBCCells(mainST_.ghosts_with_bc() -
1658  mainStencilGhost);
1659 
1660  if(executionSide_ == static_cast<int>(DomainEdgeSide::BOTH_SIDE::
1661  value)) {
1662  /* Both sides means we do not fill in directional ghost data nor need it
1663  */
1664  IntVec plus = mainStencilGhost.get_plus();
1665  plus[DirT::value] = 0;
1666  IntVec minus = mainStencilGhost.get_minus();
1667  minus[DirT::value] = 0;
1668  return GhostData(minus, plus) + mainBCCells;
1669  }
1670  else {
1671  if(executionSide_ == static_cast<int>(DomainEdgeSide::MINUS_SIDE::
1672  value)) {
1673  IntVec minus = mainStencilGhost.get_minus();
1674 
1675  minus[DirT::value] = 0;
1676 
1677  return GhostData(minus, mainStencilGhost.get_plus()) +
1678  mainBCCells;
1679  }
1680  else if(executionSide_ == static_cast<int>(DomainEdgeSide::
1681  PLUS_SIDE::value)) {
1682  IntVec plus = mainStencilGhost.get_plus();
1683 
1684  plus[DirT::value] = 0;
1685 
1686  return GhostData(mainStencilGhost.get_minus(), plus) +
1687  mainBCCells;
1688  };
1689 
1690  return mainStencilGhost + mainBCCells;
1691  };
1692  }
1693 
1694  inline GhostData varying_ghosts_without_bc(void) const {
1695  GhostData const mainStencilGhost(mainST_.ghosts_without_bc());
1696 
1697  if(executionSide_ == static_cast<int>(DomainEdgeSide::BOTH_SIDE::
1698  value)) {
1699  /* Both sides means we do not fill in directional ghost data nor need it
1700  */
1701  IntVec plus = mainStencilGhost.get_plus();
1702  plus[DirT::value] = 0;
1703  IntVec minus = mainStencilGhost.get_minus();
1704  minus[DirT::value] = 0;
1705  return GhostData(minus, plus);
1706  }
1707  else {
1708  if(executionSide_ == static_cast<int>(DomainEdgeSide::MINUS_SIDE::
1709  value)) {
1710  IntVec minus = mainStencilGhost.get_minus();
1711 
1712  minus[DirT::value] = 0;
1713 
1714  return GhostData(minus, mainStencilGhost.get_plus());
1715  }
1716  else if(executionSide_ == static_cast<int>(DomainEdgeSide::
1717  PLUS_SIDE::value)) {
1718  IntVec plus = mainStencilGhost.get_plus();
1719 
1720  plus[DirT::value] = 0;
1721 
1722  return GhostData(mainStencilGhost.get_minus(), plus);
1723  };
1724 
1725  return mainStencilGhost;
1726  };
1727  }
1728 
1730  typename StencilT::SeqWalkType,
1731  typename ConvertToSeqWalk<NegativeStencilTs>::
1732  Result,
1733  typename ConvertToSeqWalk<PositiveStencilTs>::
1734  Result,
1735  typename Arg::SeqWalkType,
1736  DirT,
1737  FieldType> typedef SeqWalkType;
1738 
1739  #ifdef ENABLE_THREADS
1741  typename StencilT::ResizeType,
1742  typename ConvertToResize<NegativeStencilTs>::
1743  Result,
1744  typename ConvertToResize<PositiveStencilTs>::
1745  Result,
1746  typename Arg::ResizeType,
1747  DirT,
1748  FieldType> typedef ResizeType;
1749  #endif
1750  /* ENABLE_THREADS */
1751 
1752  #ifdef __CUDACC__
1754  typename StencilT::GPUWalkType,
1755  typename ConvertToGPUWalk<NegativeStencilTs>::
1756  Result,
1757  typename ConvertToGPUWalk<PositiveStencilTs>::
1758  Result,
1759  typename Arg::GPUWalkType,
1760  DirT,
1761  FieldType> typedef GPUWalkType;
1762  #endif
1763  /* __CUDACC__ */
1764 
1765  NeboVaryingEdgeStencilOneDim(Arg const & a,
1766  StencilT const & mainStencil,
1767  NegativeStencilTs const & negStencils,
1768  PositiveStencilTs const & posStencils)
1769  : arg_(a), mainST_(mainStencil), negST_(negStencils), posST_(posStencils)
1770  {
1771  /* Compute which side we are executing on */
1772  {
1773  bool const minusBC = arg_.has_bc(MINUS_SIDE)[DirT::value];
1774 
1775  bool const plusBC = arg_.has_bc(PLUS_SIDE)[DirT::value];
1776 
1777  if(minusBC && plusBC) {
1778  executionSide_ = DomainEdgeSide::BOTH_SIDE::value;
1779  }
1780  else if(minusBC) {
1781  executionSide_ = DomainEdgeSide::MINUS_SIDE::value;
1782  }
1783  else if(plusBC) {
1784  executionSide_ = DomainEdgeSide::PLUS_SIDE::value;
1785  }
1786  else { executionSide_ = DomainEdgeSide::NO_SIDE::value; };
1787  };
1788 
1789  #ifndef NDEBUG
1790  GhostData const baseArgGhost(arg_.ghosts_without_bc());
1791  GhostData const neededGhost(mainST_.ghosts_without_bc() -
1792  baseArgGhost);
1793  /* Checking enough negative stencils given based on stencil width
1794  */
1795  assert(executionSide_ == static_cast<int>(DomainEdgeSide::
1796  PLUS_SIDE::value) ||
1797  executionSide_ == static_cast<int>(DomainEdgeSide::
1798  NO_SIDE::value) ||
1799  neededGhost.get_minus(DirT::value) == -(NegativeStencilTs::
1800  length));
1801  /* Checking enough positive stencils given based on stencil width
1802  */
1803  assert(executionSide_ == static_cast<int>(DomainEdgeSide::
1804  MINUS_SIDE::value) ||
1805  executionSide_ == static_cast<int>(DomainEdgeSide::
1806  NO_SIDE::value) ||
1807  neededGhost.get_plus(DirT::value) == -(PositiveStencilTs::
1808  length));
1809  /* Checking no ghost cells used on minus side if it is a domain edge
1810  */
1811  if(executionSide_ != static_cast<int>(DomainEdgeSide::PLUS_SIDE::
1812  value) && executionSide_
1813  != static_cast<int>(DomainEdgeSide::NO_SIDE::value)) {
1814  assert((VerifyNoGhostUseOnEdge<NegativeStencilTs,
1816  verify(negST_, baseArgGhost, true)));
1817  }
1818  /* Checking no ghost cells used on plus side if it is a domain edge
1819  */
1820  if(executionSide_ != static_cast<int>(DomainEdgeSide::MINUS_SIDE::
1821  value) && executionSide_
1822  != static_cast<int>(DomainEdgeSide::NO_SIDE::value)) {
1823  assert((VerifyNoGhostUseOnEdge<PositiveStencilTs,
1825  verify(posST_, baseArgGhost, true)));
1826  }
1827  #endif
1828  /* NDEBUG */;
1829  }
1830 
1831  inline GhostData ghosts_with_bc(void) const {
1832  return varying_ghosts_with_bc();
1833  }
1834 
1835  inline GhostData ghosts_without_bc(void) const {
1836  return varying_ghosts_without_bc();
1837  }
1838 
1839  inline bool has_extents(void) const { return arg_.has_extents(); }
1840 
1841  inline IntVec extents(void) const { return arg_.extents(); }
1842 
1843  inline IntVec has_bc(BCSide const bcSide) const {
1844  return arg_.has_bc(bcSide);
1845  }
1846 
1847  inline SeqWalkType init(IntVec const & extents,
1848  GhostData const & ghosts,
1849  IntVec const & hasBC,
1850  NeboOptionalArg & optArg) const {
1851  return SeqWalkType(extents + ghosts.get_plus(),
1852  arg_.init(extents, ghosts, hasBC, optArg ),
1853  mainST_.init(extents, ghosts, hasBC, optArg),
1854  ConvertToSeqWalk<NegativeStencilTs>::
1855  convert_runtime(negST_, extents, ghosts, hasBC, optArg),
1856  ConvertToSeqWalk<PositiveStencilTs>::
1857  convert_runtime(posST_, extents, ghosts, hasBC, optArg),
1858  executionSide_);
1859  }
1860 
1861  #ifdef ENABLE_THREADS
1862  inline ResizeType resize(void) const {
1863  return ResizeType(arg_.resize(),
1864  mainST_.resize(),
1865  ConvertToResize<NegativeStencilTs>::
1866  convert_runtime(negST_),
1867  ConvertToResize<PositiveStencilTs>::
1868  convert_runtime(posST_),
1869  executionSide_);
1870  }
1871  #endif
1872  /* ENABLE_THREADS */
1873 
1874  #ifdef __CUDACC__
1875  inline bool cpu_ready(void) const { return arg_.cpu_ready(); }
1876 
1877  inline bool gpu_ready(int const deviceIndex) const {
1878  return arg_.gpu_ready(deviceIndex);
1879  }
1880 
1881  inline GPUWalkType gpu_init(IntVec const & extents,
1882  GhostData const & ghosts,
1883  IntVec const & hasBC,
1884  int const deviceIndex,
1885  cudaStream_t const & lhsStream,
1886  NeboOptionalArg & optArg) const {
1887  return GPUWalkType(extents + ghosts.get_plus(),
1888  arg_.gpu_init(extents,
1889  ghosts,
1890  hasBC,
1891  deviceIndex,
1892  lhsStream,
1893  optArg),
1894  mainST_.gpu_init(extents,
1895  ghosts,
1896  hasBC,
1897  deviceIndex,
1898  lhsStream,
1899  optArg),
1900  ConvertToGPUWalk<NegativeStencilTs>::
1901  convert_runtime(negST_,
1902  extents,
1903  ghosts,
1904  hasBC,
1905  deviceIndex,
1906  lhsStream,
1907  optArg),
1908  ConvertToGPUWalk<PositiveStencilTs>::
1909  convert_runtime(posST_,
1910  extents,
1911  ghosts,
1912  hasBC,
1913  deviceIndex,
1914  lhsStream,
1915  optArg),
1916  executionSide_);
1917  }
1918 
1919  inline void stream_wait_event(cudaEvent_t const & event) const {
1920  arg_.stream_wait_event(event);
1921  }
1922 
1923  #ifdef NEBO_GPU_TEST
1924  inline void gpu_prep(int const deviceIndex) const {
1925  arg_.gpu_prep(deviceIndex);
1926  }
1927  #endif
1928  /* NEBO_GPU_TEST */
1929  #endif
1930  /* __CUDACC__ */
1931 
1932  private:
1933  Arg const arg_;
1934 
1935  StencilT const mainST_;
1936 
1937  NegativeStencilTs const negST_;
1938 
1939  PositiveStencilTs const posST_;
1940 
1941  int executionSide_;
1942  };
1943  #ifdef ENABLE_THREADS
1944  template<typename StencilT,
1945  typename NegativeStencilTs,
1946  typename PositiveStencilTs,
1947  typename Arg,
1948  typename DirT,
1949  typename FieldType>
1950  struct NeboVaryingEdgeStencilOneDim<Resize,
1951  StencilT,
1952  NegativeStencilTs,
1953  PositiveStencilTs,
1954  Arg,
1955  DirT,
1956  FieldType> {
1957  public:
1958  FieldType typedef field_type;
1959 
1960  template<typename ListT, typename Dummy = void>
1961  struct ConvertToSeqWalk {
1962  public:
1963  NeboGenericTypeList<typename ListT::First::SeqWalkType,
1964  typename ConvertToSeqWalk<typename ListT::
1965  AllButFirst>::
1966  Result> typedef Result;
1967 
1968  static inline Result convert_runtime(ListT const list,
1969  IntVec const & extents,
1970  GhostData const & ghosts,
1971  IntVec const & hasBC,
1972  NeboOptionalArg & optArg) {
1973  return ConvertToSeqWalk<typename ListT::AllButFirst>::
1974  convert_runtime(list.others(), extents, ghosts, hasBC, optArg)(list.current().init(extents,
1975  ghosts,
1976  hasBC,
1977  optArg));
1978  }
1979  };
1980  template<typename Dummy>
1981  struct ConvertToSeqWalk<NeboGenericEmptyTypeList, Dummy> {
1982  public:
1983  NeboGenericEmptyTypeList typedef Result;
1984 
1985  static inline Result convert_runtime(NeboGenericEmptyTypeList
1986  const list,
1987  IntVec const & extents,
1988  GhostData const & ghosts,
1989  IntVec const & hasBC,
1990  NeboOptionalArg & optArg) {
1991  return list;
1992  }
1993  };
1994 
1996  typename StencilT::SeqWalkType,
1997  typename ConvertToSeqWalk<NegativeStencilTs>::
1998  Result,
1999  typename ConvertToSeqWalk<PositiveStencilTs>::
2000  Result,
2001  typename Arg::SeqWalkType,
2002  DirT,
2003  FieldType> typedef SeqWalkType;
2004 
2005  NeboVaryingEdgeStencilOneDim(Arg const & a,
2006  StencilT const & mainStencil,
2007  NegativeStencilTs const & negStencils,
2008  PositiveStencilTs const & posStencils,
2009  int const executionSide)
2010  : arg_(a),
2011  mainST_(mainStencil),
2012  negST_(negStencils),
2013  posST_(posStencils),
2014  executionSide_(executionSide)
2015  {}
2016 
2017  inline SeqWalkType init(IntVec const & extents,
2018  GhostData const & ghosts,
2019  IntVec const & hasBC,
2020  NeboOptionalArg & optArg) const {
2021  return SeqWalkType(extents + ghosts.get_plus(),
2022  arg_.init(extents, ghosts, hasBC, optArg),
2023  mainST_.init(extents, ghosts, hasBC, optArg),
2024  ConvertToSeqWalk<NegativeStencilTs>::
2025  convert_runtime(negST_,
2026  extents,
2027  ghosts,
2028  hasBC,
2029  optArg),
2030  ConvertToSeqWalk<PositiveStencilTs>::
2031  convert_runtime(posST_,
2032  extents,
2033  ghosts,
2034  hasBC,
2035  optArg),
2036  executionSide_);
2037  }
2038 
2039  private:
2040  Arg const arg_;
2041 
2042  StencilT const mainST_;
2043 
2044  NegativeStencilTs const negST_;
2045 
2046  PositiveStencilTs const posST_;
2047 
2048  int const executionSide_;
2049  }
2050  #endif
2051  /* ENABLE_THREADS */;
2052  template<typename StencilT,
2053  typename NegativeStencilTs,
2054  typename PositiveStencilTs,
2055  typename Arg,
2056  typename DirT,
2057  typename FieldType>
2059  StencilT,
2060  NegativeStencilTs,
2061  PositiveStencilTs,
2062  Arg,
2063  DirT,
2064  FieldType> {
2065  public:
2066  FieldType typedef field_type;
2067 
2068  typename field_type::value_type typedef value_type;
2069 
2071  StencilT,
2072  NegativeStencilTs,
2073  PositiveStencilTs,
2074  Arg,
2075  DirT,
2076  FieldType> typedef MyType;
2077 
2078  template<typename ListT, int CurrentIndex>
2079  struct EvalEdgeStencil {
2080  template<typename OptionalArgT>
2081  static inline value_type eval(ListT const & list,
2082  int const searchIndex,
2083  int const x,
2084  int const y,
2085  int const z) {
2086  return (CurrentIndex == searchIndex ? list.current().template
2087  eval<OptionalArgT>(x,
2088  y,
2089  z)
2090  : EvalEdgeStencil<typename ListT::Collection,
2091  CurrentIndex - 1>::template eval<OptionalArgT>(list.others(),
2092  searchIndex,
2093  x,
2094  y,
2095  z));
2096  }
2097  };
2098  template<typename ListT>
2099  struct EvalEdgeStencil<ListT, 0> {
2100  template<typename OptionalArgT>
2101  static inline value_type eval(ListT const & list,
2102  int const searchIndex,
2103  int const x,
2104  int const y,
2105  int const z) {
2106  return list.current().template eval<OptionalArgT>(x, y, z);
2107  }
2108  };
2109 
2110  template<typename IndexDirT, typename Dummy = void>
2111  struct GetIndex;
2112  template<typename Dummy>
2113  struct GetIndex<XDIR, Dummy> {
2114  public:
2115  static inline int index(int const x, int const y, int const z) {
2116  return x;
2117  }
2118  };
2119  template<typename Dummy>
2120  struct GetIndex<YDIR, Dummy> {
2121  public:
2122  static inline int index(int const x, int const y, int const z) {
2123  return y;
2124  }
2125  };
2126  template<typename Dummy>
2127  struct GetIndex<ZDIR, Dummy> {
2128  public:
2129  static inline int index(int const x, int const y, int const z) {
2130  return z;
2131  }
2132  };
2133 
2134  template<typename T, typename Dummy = void>
2135  struct SpecialArgHandler {
2136  constexpr static inline bool hasBC(void) { return false; }
2137  };
2138  template<typename Dummy>
2139  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::HasBCOnX,
2140  Dummy> {
2141  constexpr static inline bool hasBC(void) {
2142  return std::is_same<DirT, XDIR>::value;
2143  }
2144  };
2145  template<typename Dummy>
2146  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::HasBCOnY,
2147  Dummy> {
2148  constexpr static inline bool hasBC(void) {
2149  return std::is_same<DirT, YDIR>::value;
2150  }
2151  };
2152  template<typename Dummy>
2153  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::HasBCOnZ,
2154  Dummy> {
2155  constexpr static inline bool hasBC(void) {
2156  return std::is_same<DirT, ZDIR>::value;
2157  }
2158  };
2159  template<typename ... OptionalArgT>
2160  struct OptionalArgParser;
2161  template<typename ... T>
2162  struct OptionalArgParser<CompileTimeOptionalArgs<T ...> > {
2163  constexpr static inline bool hasBC(void) {
2164  return OptionalArgParser<T ...>::hasBC();
2165  }
2166  };
2167  template<typename U, typename ... T>
2168  struct OptionalArgParser<U, T ...> {
2169  constexpr static inline bool hasBC(void) {
2170  return SpecialArgHandler<U>::hasBC() || OptionalArgParser<T ...>::
2171  hasBC();
2172  }
2173  };
2174  template<typename T>
2175  struct OptionalArgParser<T> {
2176  constexpr static inline bool hasBC(void) {
2177  return SpecialArgHandler<T>::hasBC() || false;
2178  }
2179  };
2180 
2181  template<typename OptionalArgT>
2182  inline value_type evalOnSide(int const x, int const y, int const z) const {
2183  if(OptionalArgParser<OptionalArgT>::hasBC()) {
2184  if(executionSide_ != static_cast<int>(DomainEdgeSide::NO_SIDE::
2185  value)) {
2186  int const i = GetIndex<DirT>::index(x, y, z);
2187 
2188  int const negi = (dirExtent_ - 1) - i;
2189 
2190  if(executionSide_ != static_cast<int>(DomainEdgeSide::
2191  MINUS_SIDE::value) &&
2192  0 <= negi && negi < PositiveStencilTs::length) {
2193  return EvalEdgeStencil<PositiveStencilTs,
2194  PositiveStencilTs::length - 1>::
2195  template eval<OptionalArgT>(posST_, negi, x, y, z);
2196  }
2197  else if(executionSide_ != static_cast<int>(DomainEdgeSide::
2198  PLUS_SIDE::value)
2199  && 0 <= i && i < NegativeStencilTs::length) {
2200  return EvalEdgeStencil<NegativeStencilTs,
2201  NegativeStencilTs::length - 1>::
2202  template eval<OptionalArgT>(negST_, i, x, y, z);
2203  };
2204  };
2205  };
2206 
2207  return mainST_.template eval<OptionalArgT>(x, y, z);
2208  }
2209 
2210  NeboVaryingEdgeStencilOneDim(IntVec const & plusLimit,
2211  Arg const & arg,
2212  StencilT const & mainStencil,
2213  NegativeStencilTs const & negStencils,
2214  PositiveStencilTs const & posStencils,
2215  int const executionSide)
2216  : dirExtent_(GetIndex<DirT>::index(plusLimit[0],
2217  plusLimit[1],
2218  plusLimit[2])),
2219  mainST_(mainStencil),
2220  negST_(negStencils),
2221  posST_(posStencils),
2222  executionSide_(executionSide)
2223  {}
2224 
2225  template<typename OptionalArgT>
2226  inline value_type eval(int const x, int const y, int const z) const {
2227  return evalOnSide<OptionalArgT>(x, y, z);
2228  }
2229 
2230  private:
2231  int const dirExtent_;
2232 
2233  StencilT const mainST_;
2234 
2235  NegativeStencilTs const negST_;
2236 
2237  PositiveStencilTs const posST_;
2238 
2239  int const executionSide_;
2240  };
2241  #ifdef __CUDACC__
2242  template<typename StencilT,
2243  typename NegativeStencilTs,
2244  typename PositiveStencilTs,
2245  typename Arg,
2246  typename DirT,
2247  typename FieldType>
2248  struct NeboVaryingEdgeStencilOneDim<GPUWalk,
2249  StencilT,
2250  NegativeStencilTs,
2251  PositiveStencilTs,
2252  Arg,
2253  DirT,
2254  FieldType> {
2255  public:
2256  FieldType typedef field_type;
2257 
2258  typename field_type::value_type typedef value_type;
2259 
2260  template<typename ListT, int CurrentIndex>
2261  struct EvalEdgeStencil {
2262  template<typename OptionalArgT>
2263  __device__ static inline value_type eval(ListT const & list,
2264  int const searchIndex,
2265  int const x,
2266  int const y,
2267  int const z) {
2268  return (CurrentIndex == searchIndex ? list.current().template
2269  eval<OptionalArgT>(x,
2270  y,
2271  z)
2272  : EvalEdgeStencil<typename ListT::Collection,
2273  CurrentIndex - 1>::template eval<OptionalArgT>(list.others(),
2274  searchIndex,
2275  x,
2276  y,
2277  z));
2278  }
2279  };
2280  template<typename ListT>
2281  struct EvalEdgeStencil<ListT, 0> {
2282  template<typename OptionalArgT>
2283  __device__ static inline value_type eval(ListT const & list,
2284  int const searchIndex,
2285  int const x,
2286  int const y,
2287  int const z) {
2288  return list.current().template eval<OptionalArgT>(x, y, z);
2289  }
2290  };
2291 
2292  template<typename IndexDirT, typename Dummy = void>
2293  struct GetIndex;
2294  template<typename Dummy>
2295  struct GetIndex<XDIR, Dummy> {
2296  public:
2297  #ifdef __CUDACC__
2298  __host__ __device__
2299  #endif
2300  /* __CUDACC__ */ static inline int index(int const x,
2301  int const y,
2302  int const z) {
2303  return x;
2304  }
2305  };
2306  template<typename Dummy>
2307  struct GetIndex<YDIR, Dummy> {
2308  public:
2309  #ifdef __CUDACC__
2310  __host__ __device__
2311  #endif
2312  /* __CUDACC__ */ static inline int index(int const x,
2313  int const y,
2314  int const z) {
2315  return y;
2316  }
2317  };
2318  template<typename Dummy>
2319  struct GetIndex<ZDIR, Dummy> {
2320  public:
2321  #ifdef __CUDACC__
2322  __host__ __device__
2323  #endif
2324  /* __CUDACC__ */ static inline int index(int const x,
2325  int const y,
2326  int const z) {
2327  return z;
2328  }
2329  };
2330 
2331  template<typename T, typename Dummy = void>
2332  struct SpecialArgHandler {
2333  __device__ constexpr static inline bool hasBC(void) {
2334  return false;
2335  }
2336  };
2337  template<typename Dummy>
2338  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::
2339  HasBCOnX,
2340  Dummy> {
2341  __device__ constexpr static inline bool hasBC(void) {
2342  return std::is_same<DirT, XDIR>::value;
2343  }
2344  };
2345  template<typename Dummy>
2346  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::
2347  HasBCOnY,
2348  Dummy> {
2349  __device__ constexpr static inline bool hasBC(void) {
2350  return std::is_same<DirT, YDIR>::value;
2351  }
2352  };
2353  template<typename Dummy>
2354  struct SpecialArgHandler<CompileTimeOptionalArgsNamespace::
2355  HasBCOnZ,
2356  Dummy> {
2357  __device__ constexpr static inline bool hasBC(void) {
2358  return std::is_same<DirT, ZDIR>::value;
2359  }
2360  };
2361  template<typename ... OptionalArgT>
2362  struct OptionalArgParser;
2363  template<typename ... T>
2364  struct OptionalArgParser<CompileTimeOptionalArgs<T ...> > {
2365  __device__ constexpr static inline bool hasBC(void) {
2366  return OptionalArgParser<T ...>::hasBC();
2367  }
2368  };
2369  template<typename U, typename ... T>
2370  struct OptionalArgParser<U, T ...> {
2371  __device__ constexpr static inline bool hasBC(void) {
2372  return SpecialArgHandler<U>::hasBC() || OptionalArgParser<T
2373  ...>::
2374  hasBC();
2375  }
2376  };
2377  template<typename T>
2378  struct OptionalArgParser<T> {
2379  __device__ constexpr static inline bool hasBC(void) {
2380  return SpecialArgHandler<T>::hasBC() || false;
2381  }
2382  };
2383 
2384  template<typename OptionalArgT>
2385  __device__ inline value_type evalOnSide(int const x,
2386  int const y,
2387  int const z) const {
2388  if(OptionalArgParser<OptionalArgT>::hasBC()) {
2389  if(executionSide_ != static_cast<int>(DomainEdgeSide::
2390  NO_SIDE::value)) {
2391  int const i = GetIndex<DirT>::index(x, y, z);
2392 
2393  int const negi = (dirExtent_ - 1) - i;
2394 
2395  if(executionSide_ != static_cast<int>(DomainEdgeSide::
2396  MINUS_SIDE::value)
2397  && 0 <= negi && negi < PositiveStencilTs::length) {
2398  return EvalEdgeStencil<PositiveStencilTs,
2399  PositiveStencilTs::length - 1>::
2400  template eval<OptionalArgT>(posST_,
2401  negi,
2402  x,
2403  y,
2404  z);
2405  }
2406  else if(executionSide_ != static_cast<int>(DomainEdgeSide::
2407  PLUS_SIDE::
2408  value) && 0 <=
2409  i && i < NegativeStencilTs::length) {
2410  return EvalEdgeStencil<NegativeStencilTs,
2411  NegativeStencilTs::length - 1>::
2412  template eval<OptionalArgT>(negST_, i, x, y, z);
2413  };
2414  };
2415  };
2416 
2417  return mainST_.template eval<OptionalArgT>(x, y, z);
2418  }
2419 
2420  NeboVaryingEdgeStencilOneDim(IntVec const & extents,
2421  Arg const & arg,
2422  StencilT const & mainStencil,
2423  NegativeStencilTs const & negStencils,
2424  PositiveStencilTs const & posStencils,
2425  int const executionSide)
2426  : dirExtent_(GetIndex<DirT>::index(extents[0],
2427  extents[1],
2428  extents[2])),
2429  mainST_(mainStencil),
2430  negST_(negStencils),
2431  posST_(posStencils),
2432  executionSide_(executionSide)
2433  {}
2434 
2435  template<typename OptionalArgT>
2436  __device__ inline value_type eval(int const x,
2437  int const y,
2438  int const z) const {
2439  return evalOnSide<OptionalArgT>(x, y, z);
2440  }
2441 
2442  private:
2443  int const dirExtent_;
2444 
2445  StencilT const mainST_;
2446 
2447  NegativeStencilTs const negST_;
2448 
2449  PositiveStencilTs const posST_;
2450 
2451  int const executionSide_;
2452  }
2453  #endif
2454  /* __CUDACC__ */;
2455  } /* SpatialOps */
2456 
2457 #endif
2458 /* 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
Parameter used to initialize Nebo expression operands across modes. The argument only stores informat...
Definition: NeboBasic.h:312
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...