SpatialOps
NeboRhs.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_RHS_H
26  #define NEBO_RHS_H
27 
28  namespace SpatialOps {
29  template<typename CurrentMode, typename AtomicType>
30  struct NeboScalar;
31  template<typename AtomicType>
32  struct NeboScalar<Initial, AtomicType> {
33  public:
34  AtomicType typedef value_type;
35 
36  NeboScalar<SeqWalk, AtomicType> typedef SeqWalkType;
37 
38  #ifdef ENABLE_THREADS
39  NeboScalar<Resize, AtomicType> typedef ResizeType;
40  #endif
41  /* ENABLE_THREADS */
42 
43  #ifdef __CUDACC__
44  NeboScalar<GPUWalk, AtomicType> typedef GPUWalkType;
45  #endif
46  /* __CUDACC__ */
47 
48  NeboScalar(value_type const v)
49  : value_(v)
50  {}
51 
52  inline GhostData ghosts_with_bc(void) const {
53  return GhostData(GHOST_MAX);
54  }
55 
56  inline GhostData ghosts_without_bc(void) const {
57  return GhostData(GHOST_MAX);
58  }
59 
60  inline bool has_extents(void) const { return false; }
61 
62  inline IntVec extents(void) const { return IntVec(0, 0, 0); }
63 
64  inline IntVec has_bc(BCSide const bcSide) const {
65  return IntVec(0, 0, 0);
66  }
67 
68  inline SeqWalkType init(IntVec const & extents,
69  GhostData const & ghosts,
70  IntVec const & hasBC,
71  NeboOptionalArg const & optArg ) const {
72  return SeqWalkType(value_);
73  }
74 
75  #ifdef ENABLE_THREADS
76  inline ResizeType resize(void) const { return ResizeType(value_); }
77  #endif
78  /* ENABLE_THREADS */
79 
80  #ifdef __CUDACC__
81  inline bool cpu_ready(void) const { return true; }
82 
83  inline bool gpu_ready(int const deviceIndex) const { return true; }
84 
85  inline GPUWalkType gpu_init(IntVec const & extents,
86  GhostData const & ghosts,
87  IntVec const & hasBC,
88  int const deviceIndex,
89  cudaStream_t const & lhsStream,
90  NeboOptionalArg & optArg) const {
91  return GPUWalkType(value_);
92  }
93 
94  inline void stream_wait_event(cudaEvent_t const & event) const {}
95 
96  #ifdef NEBO_GPU_TEST
97  inline void gpu_prep(int const deviceIndex) const {}
98  #endif
99  /* NEBO_GPU_TEST */
100  #endif
101  /* __CUDACC__ */
102 
103  private:
104  value_type const value_;
105  };
106  #ifdef ENABLE_THREADS
107  template<typename AtomicType>
108  struct NeboScalar<Resize, AtomicType> {
109  public:
110  AtomicType typedef value_type;
111 
112  NeboScalar<SeqWalk, AtomicType> typedef SeqWalkType;
113 
114  NeboScalar(value_type const value)
115  : value_(value)
116  {}
117 
118  inline SeqWalkType init(IntVec const & extents,
119  GhostData const & ghosts,
120  IntVec const & hasBC,
121  NeboOptionalArg & optArg) const {
122  return SeqWalkType(value_);
123  }
124 
125  private:
126  value_type const value_;
127  }
128  #endif
129  /* ENABLE_THREADS */;
130  template<typename AtomicType>
131  struct NeboScalar<SeqWalk, AtomicType> {
132  public:
133  AtomicType typedef value_type;
134 
135  NeboScalar(value_type const value)
136  : value_(value)
137  {}
138 
139  template<typename OptionalArgT>
140  inline value_type eval(int const x, int const y, int const z) const {
141  return value_;
142  }
143 
144  private:
145  value_type const value_;
146  };
147  #ifdef __CUDACC__
148  template<typename AtomicType>
149  struct NeboScalar<GPUWalk, AtomicType> {
150  public:
151  AtomicType typedef value_type;
152 
153  NeboScalar(value_type const value)
154  : value_(value)
155  {}
156 
157  template<typename OptionalArgT>
158  __device__ inline value_type eval(int const x,
159  int const y,
160  int const z) const {
161  return value_;
162  }
163 
164  private:
165  value_type const value_;
166  }
167  #endif
168  /* __CUDACC__ */;
169 
170  template<typename CurrentMode, typename FieldType>
172  template<typename FieldType>
173  struct NeboConstField<Initial, FieldType> {
174  public:
175  FieldType typedef field_type;
176 
177  NeboConstField<SeqWalk, FieldType> typedef SeqWalkType;
178 
179  #ifdef ENABLE_THREADS
180  NeboConstField<Resize, FieldType> typedef ResizeType;
181  #endif
182  /* ENABLE_THREADS */
183 
184  #ifdef __CUDACC__
185  NeboConstField<GPUWalk, FieldType> typedef GPUWalkType;
186  #endif
187  /* __CUDACC__ */
188 
189  NeboConstField(FieldType const & f)
190  : field_(f)
191  {}
192 
193  inline GhostData ghosts_with_bc(void) const {
194  return field_.get_valid_ghost_data() + point_to_ghost(field_.boundary_info().has_extra());
195  }
196 
197  inline GhostData ghosts_without_bc(void) const {
198  return field_.get_valid_ghost_data();
199  }
200 
201  inline bool has_extents(void) const { return true; }
202 
203  inline IntVec extents(void) const {
204  return field_.window_with_ghost().extent() - field_.get_valid_ghost_data().get_minus()
205  - field_.get_valid_ghost_data().get_plus();
206  }
207 
208  inline IntVec has_bc(BCSide const bcSide) const {
209  return field_.boundary_info().has_bc(bcSide);
210  }
211 
212  inline SeqWalkType init(IntVec const & extents,
213  GhostData const & ghosts,
214  IntVec const & hasBC,
215  NeboOptionalArg const & optArg) const {
216  return SeqWalkType(field_);
217  }
218 
219  #ifdef ENABLE_THREADS
220  inline ResizeType resize(void) const { return ResizeType(field_); }
221  #endif
222  /* ENABLE_THREADS */
223 
224  #ifdef __CUDACC__
225  inline bool cpu_ready(void) const {
226  return field_.is_valid(CPU_INDEX);
227  }
228 
229  inline bool gpu_ready(int const deviceIndex) const {
230  return field_.is_valid(deviceIndex);
231  }
232 
233  inline GPUWalkType gpu_init(IntVec const & extents,
234  GhostData const & ghosts,
235  IntVec const & hasBC,
236  int const deviceIndex,
237  cudaStream_t const & lhsStream,
238  NeboOptionalArg & optArg) const {
239  return GPUWalkType(lhsStream, deviceIndex, field_);
240  }
241 
242  inline void stream_wait_event(cudaEvent_t const & event) const {
243  cudaStreamWaitEvent(field_.get_stream(), event, 0);
244  }
245 
246  #ifdef NEBO_GPU_TEST
247  inline void gpu_prep(int const deviceIndex) const {
248  const_cast<FieldType *>(&field_)->add_device(deviceIndex);
249  }
250  #endif
251  /* NEBO_GPU_TEST */
252  #endif
253  /* __CUDACC__ */
254 
255  private:
256  FieldType const field_;
257  };
258  #ifdef ENABLE_THREADS
259  template<typename FieldType>
260  struct NeboConstField<Resize, FieldType> {
261  public:
262  FieldType typedef field_type;
263 
264  NeboConstField<SeqWalk, FieldType> typedef SeqWalkType;
265 
266  NeboConstField(FieldType const & f)
267  : field_(f)
268  {}
269 
270  inline SeqWalkType init(IntVec const & extents,
271  GhostData const & ghosts,
272  IntVec const & hasBC,
273  NeboOptionalArg & optArg) const {
274  return SeqWalkType(field_);
275  }
276 
277  private:
278  FieldType const field_;
279  }
280  #endif
281  /* ENABLE_THREADS */;
282  template<typename FieldType>
283  struct NeboConstField<SeqWalk, FieldType> {
284  public:
285  FieldType typedef field_type;
286 
287  typename field_type::value_type typedef value_type;
288 
289  NeboConstField(FieldType const & f)
290  : xGlob_(f.window_with_ghost().glob_dim(0)),
291  yGlob_(f.window_with_ghost().glob_dim(1)),
292  base_(f.field_values(CPU_INDEX) + (f.window_with_ghost().offset(0) +
293  f.get_valid_ghost_data().get_minus(0))
294  + (f.window_with_ghost().glob_dim(0) * ((f.window_with_ghost().offset(1)
295  + f.get_valid_ghost_data().get_minus(1))
296  + (f.window_with_ghost().glob_dim(1)
297  * (f.window_with_ghost().offset(2)
298  + f.get_valid_ghost_data().get_minus(2))))))
299 
300  {}
301 
302  template<typename OptionalArgT>
303  inline value_type eval(int const x, int const y, int const z) const {
304  return base_[x + xGlob_ * (y + (yGlob_ * z))];
305  }
306 
307  private:
308  int const xGlob_;
309 
310  int const yGlob_;
311 
312  value_type const * base_;
313 
314  };
315  #ifdef __CUDACC__
316  template<typename FieldType>
317  struct NeboConstField<GPUWalk, FieldType> {
318  public:
319  FieldType typedef field_type;
320 
321  typename field_type::value_type typedef value_type;
322 
323  NeboConstField(cudaStream_t const & lhsStream,
324  int const deviceIndex,
325  FieldType const & f)
326  : base_(f.field_values(deviceIndex) + (f.window_with_ghost().offset(0)
327  + f.get_valid_ghost_data().get_minus(0))
328  + (f.window_with_ghost().glob_dim(0) * ((f.window_with_ghost().offset(1)
329  + f.get_valid_ghost_data().get_minus(1))
330  + (f.window_with_ghost().glob_dim(1)
331  * (f.window_with_ghost().offset(2)
332  + f.get_valid_ghost_data().get_minus(2)))))),
333  xGlob_(f.window_with_ghost().glob_dim(0)),
334  yGlob_(f.window_with_ghost().glob_dim(1))
335  { cudaStreamWaitEvent(lhsStream, f.get_last_event(), 0); }
336 
337  template<typename OptionalArgT>
338  __device__ inline value_type eval(int const x,
339  int const y,
340  int const z) const {
341  return base_[x + xGlob_ * (y + (yGlob_ * z))];
342  }
343 
344  private:
345  value_type const * base_;
346 
347  int const xGlob_;
348 
349  int const yGlob_;
350  }
351  #endif
352  /* __CUDACC__ */;
353 
354  template<typename CurrentMode, typename T>
356  template<typename T>
357  struct NeboConstSingleValueField<Initial, T> {
358  public:
360  field_type;
361 
363  SingleValueFieldType;
364 
365  NeboConstSingleValueField<SeqWalk, T> typedef SeqWalkType;
366 
367  #ifdef ENABLE_THREADS
368  NeboConstSingleValueField<Resize, T> typedef ResizeType;
369  #endif
370  /* ENABLE_THREADS */
371 
372  #ifdef __CUDACC__
373  NeboConstSingleValueField<GPUWalk, T> typedef GPUWalkType;
374  #endif
375  /* __CUDACC__ */
376 
377  NeboConstSingleValueField(SingleValueFieldType const & f)
378  : field_(f)
379  {}
380 
381  inline GhostData ghosts_with_bc(void) const {
382  return GhostData(GHOST_MAX);
383  }
384 
385  inline GhostData ghosts_without_bc(void) const {
386  return GhostData(GHOST_MAX);
387  }
388 
389  inline bool has_extents(void) const { return false; }
390 
391  inline IntVec extents(void) const { return IntVec(0, 0, 0); }
392 
393  inline IntVec has_bc(BCSide const bcSide) const {
394  return IntVec(0, 0, 0);
395  }
396 
397  inline SeqWalkType init(IntVec const & extents,
398  GhostData const & ghosts,
399  IntVec const & hasBC,
400  NeboOptionalArg & optArg) const {
401  return SeqWalkType(* field_.field_values(CPU_INDEX));
402  }
403 
404  #ifdef ENABLE_THREADS
405  inline ResizeType resize(void) const {
406  return ResizeType(* field_.field_values(CPU_INDEX));
407  }
408  #endif
409  /* ENABLE_THREADS */
410 
411  #ifdef __CUDACC__
412  inline bool cpu_ready(void) const {
413  return field_.is_valid(CPU_INDEX);
414  }
415 
416  inline bool gpu_ready(int const deviceIndex) const {
417  return field_.is_valid(deviceIndex);
418  }
419 
420  inline GPUWalkType gpu_init(IntVec const & extents,
421  GhostData const & ghosts,
422  IntVec const & hasBC,
423  int const deviceIndex,
424  cudaStream_t const & lhsStream,
425  NeboOptionalArg & optArg) const {
426  return GPUWalkType(lhsStream, deviceIndex, field_);
427  }
428 
429  inline void stream_wait_event(cudaEvent_t const & event) const {
430  cudaStreamWaitEvent(field_.get_stream(), event, 0);
431  }
432 
433  #ifdef NEBO_GPU_TEST
434  inline void gpu_prep(int const deviceIndex) const {
435  const_cast<SingleValueFieldType *>(&field_)->add_device(deviceIndex);
436  }
437  #endif
438  /* NEBO_GPU_TEST */
439  #endif
440  /* __CUDACC__ */
441 
442  private:
443  SingleValueFieldType const field_;
444  };
445  #ifdef ENABLE_THREADS
446  template<typename T>
447  struct NeboConstSingleValueField<Resize, T> {
448  public:
450  field_type;
451 
452  NeboConstSingleValueField<SeqWalk, T> typedef SeqWalkType;
453 
454  NeboConstSingleValueField(double const & v)
455  : value_(v)
456  {}
457 
458  inline SeqWalkType init(IntVec const & extents,
459  GhostData const & ghosts,
460  IntVec const & hasBC,
461  NeboOptionalArg & optArg) const {
462  return SeqWalkType(value_);
463  }
464 
465  private:
466  double const value_;
467  }
468  #endif
469  /* ENABLE_THREADS */;
470  template<typename T>
471  struct NeboConstSingleValueField<SeqWalk, T> {
472  public:
474  field_type;
475 
476  typename field_type::value_type typedef value_type;
477 
478  NeboConstSingleValueField(double const & v)
479  : value_(v)
480  {}
481 
482  template<typename OptionalArgT>
483  inline value_type eval(int const x, int const y, int const z) const {
484  return value_;
485  }
486 
487  private:
488  double value_;
489  };
490  #ifdef __CUDACC__
491  template<typename T>
492  struct NeboConstSingleValueField<GPUWalk, T> {
493  public:
495  field_type;
496 
497  typename field_type::value_type typedef value_type;
498 
500  SingleValueFieldType;
501 
502  NeboConstSingleValueField(cudaStream_t const & lhsStream,
503  int const deviceIndex,
504  SingleValueFieldType const & f)
505  : pointer_(f.field_values(deviceIndex))
506  { cudaStreamWaitEvent(lhsStream, f.get_last_event(), 0); }
507 
508  template<typename OptionalArgT>
509  __device__ inline value_type eval(int const x,
510  int const y,
511  int const z) const {
512  return *pointer_;
513  }
514 
515  private:
516  value_type const * pointer_;
517  }
518  #endif
519  /* __CUDACC__ */;
520 
521  template<typename CurrentMode, typename MapperState>
523  template<typename MapperState>
524  struct NeboIndexMapper<Initial, MapperState> {
525  public:
526  MapperState typedef mapper_state;
527 
528  NeboIndexMapper<SeqWalk, MapperState> typedef SeqWalkType;
529 
530  #ifdef ENABLE_THREADS
531  NeboIndexMapper<Resize, MapperState> typedef ResizeType;
532  #endif
533  /* ENABLE_THREADS */
534 
535  #ifdef __CUDACC__
536  NeboIndexMapper<GPUWalk, MapperState> typedef GPUWalkType;
537  #endif
538  /* __CUDACC__ */
539 
540  NeboIndexMapper(const MapperState & m)
541  : mapperState_(m)
542  {}
543 
544  inline SeqWalkType init(NeboOptionalArg & optArg) const {
545  return SeqWalkType(mapperState_, optArg.mappedValueRef());
546  }
547 
548  #ifdef ENABLE_THREADS
549  inline ResizeType resize() const { return ResizeType(mapperState_); }
550  #endif
551  /* ENABLE_THREADS */
552 
553  #ifdef __CUDACC__
554  inline bool cpu_ready(void) const {
555  return true;
556  }
557 
558  inline bool gpu_ready(int const deviceIndex) const {
559  return true;
560  }
561 
562  inline GPUWalkType gpu_init(int const deviceIndex,
563  cudaStream_t const & lhsStream,
564  NeboOptionalArg & optArg) const {
565  return GPUWalkType(lhsStream, deviceIndex, mapperState_, optArg.mappedValueRef());
566  }
567 
568  inline void stream_wait_event(cudaEvent_t const & event) const {
569 
570  }
571 
572  #ifdef NEBO_GPU_TEST
573  inline void gpu_prep(int const deviceIndex) const {
574 
575  }
576  #endif
577  /* NEBO_GPU_TEST */
578  #endif
579  /* __CUDACC__ */
580 
581  private:
582  MapperState mapperState_;
583  };
584  #ifdef ENABLE_THREADS
585  template<typename MapperState>
586  struct NeboIndexMapper<Resize, MapperState> {
587  public:
588  MapperState typedef mapper_state;
589 
590  NeboIndexMapper<SeqWalk, MapperState> typedef SeqWalkType;
591 
592  NeboIndexMapper(const MapperState & m)
593  : mapperState_(m)
594  {}
595 
596  inline SeqWalkType init(NeboOptionalArg & optArg) const {
597  return SeqWalkType(mapperState_, optArg.mappedValueRef());
598  }
599 
600  private:
601  MapperState mapperState_;
602  }
603  #endif
604  /* ENABLE_THREADS */;
605  template<typename MapperState>
606  struct NeboIndexMapper<SeqWalk, MapperState> {
607  public:
608  MapperState typedef mapper_state;
609  typename MapperState::p_iterator typedef iterator;
610 
611  NeboIndexMapper(const mapper_state & m, double * mappedValue)
612  : mapperState_(m), mappedValue_(mappedValue)
613  {}
614 
615  inline double* mappedValueRef (void) { return mappedValue_; }
616 
617  inline void eval(int const x, int const y, int const z, iterator it) {
618 
619  return mapperState_.eval(x,y,z,it,mappedValue_);
620 
621  }
622 
623  inline void next (int const x, int const y, int const z, iterator prev) {
624 
625  return mapperState_.next(x,y,z,prev,mappedValue_);
626  }
627 
628 
629  private:
630 
631  MapperState mapperState_;
632  double * mappedValue_;
633  };
634  #ifdef __CUDACC__
635  template<typename MapperState>
636  struct NeboIndexMapper<GPUWalk, MapperState> {
637  public:
638  MapperState typedef mapper_state;
639  typename MapperState::p_iterator typedef iterator;
640 
641  NeboIndexMapper(cudaStream_t const & lhsStream,
642  int const deviceIndex,
643  const mapper_state & m,
644  double * mappedValue)
645  : mapperState_(m), mappedValue_(mappedValue)
646  {}
647 
648  __device__ inline double* mappedValueRef (void) { return mappedValue_; }
649 
650  __device__ inline void eval(int const x, int const y, int const z, iterator it) {
651  return mapperState_.gpu_eval(x,y,z,it,mappedValue_);
652 
653  }
654 
655  __device__ inline void next (int const x, int const y, int const z, iterator prev) {
656  return mapperState_.gpu_next(x,y,z,prev,mappedValue_);
657  }
658 
659  private:
660 
661  MapperState mapperState_;
662  double * mappedValue_;
663 
664  }
665  #endif
666  /* __CUDACC__ */;
667 
668  } /* SpatialOps */
669 
670 #endif
671 /* NEBO_RHS_H */
672 
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...