SpatialOps
NeboMask.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_MASK_H
26  #define NEBO_MASK_H
27 
28  namespace SpatialOps {
29  template<typename CurrentMode, typename FieldType>
30  struct NeboMask;
31  template<typename FieldType>
32  struct NeboMask<Initial, FieldType> {
33  public:
34  FieldType typedef field_type;
35 
36  NeboMask<SeqWalk, FieldType> typedef SeqWalkType;
37 
38  #ifdef ENABLE_THREADS
39  NeboMask<Resize, FieldType> typedef ResizeType;
40  #endif
41  /* ENABLE_THREADS */
42 
43  #ifdef __CUDACC__
44  NeboMask<GPUWalk, FieldType> typedef GPUWalkType;
45  #endif
46  /* __CUDACC__ */
47 
49  : mask_(m)
50  {}
51 
52  inline GhostData ghosts_with_bc(void) const {
53  return mask_.get_valid_ghost_data() + point_to_ghost(mask_.boundary_info().has_extra());
54  }
55 
56  inline GhostData ghosts_without_bc(void) const {
57  return mask_.get_valid_ghost_data();
58  }
59 
60  inline bool has_extents(void) const { return true; }
61 
62  inline IntVec extents(void) const {
63  return mask_.window_without_ghost().extent();
64  }
65 
66  inline IntVec has_bc(BCSide const bcSide) const {
67  return point_to_ghost(mask_.boundary_info().has_bc(bcSide));
68  }
69 
70  inline SeqWalkType init(IntVec const & extents,
71  GhostData const & ghosts,
72  IntVec const & hasBC) const {
73  return SeqWalkType(mask_);
74  }
75 
76  #ifdef ENABLE_THREADS
77  inline ResizeType resize(void) const { return ResizeType(mask_); }
78  #endif
79  /* ENABLE_THREADS */
80 
81  #ifdef __CUDACC__
82  inline bool cpu_ready(void) const {
83  return mask_.find_consumer(CPU_INDEX);
84  }
85 
86  inline bool gpu_ready(int const deviceIndex) const {
87  return mask_.find_consumer(deviceIndex);
88  }
89 
90  inline GPUWalkType gpu_init(IntVec const & extents,
91  GhostData const & ghosts,
92  IntVec const & hasBC,
93  int const deviceIndex,
94  cudaStream_t const & lhsStream) const {
95  return GPUWalkType(deviceIndex, mask_);
96  }
97 
98  inline void stream_wait_event(cudaEvent_t const & event) const {}
99 
100  #ifdef NEBO_GPU_TEST
101  inline void gpu_prep(int const deviceIndex) const {
102  const_cast<SpatialMask<FieldType> *>(&mask_)->add_consumer(deviceIndex);
103  }
104  #endif
105  /* NEBO_GPU_TEST */
106  #endif
107  /* __CUDACC__ */
108 
109  private:
110  SpatialMask<FieldType> const mask_;
111  };
112  #ifdef ENABLE_THREADS
113  template<typename FieldType>
114  struct NeboMask<Resize, FieldType> {
115  public:
116  FieldType typedef field_type;
117 
118  NeboMask<SeqWalk, FieldType> typedef SeqWalkType;
119 
121  : mask_(m)
122  {}
123 
124  inline SeqWalkType init(IntVec const & extents,
125  GhostData const & ghosts,
126  IntVec const & hasBC) const {
127  return SeqWalkType(mask_);
128  }
129 
130  private:
131  SpatialMask<FieldType> const mask_;
132  }
133  #endif
134  /* ENABLE_THREADS */;
135  template<typename FieldType>
136  struct NeboMask<SeqWalk, FieldType> {
137  public:
138  FieldType typedef field_type;
139 
140  typename field_type::value_type typedef value_type;
141 
143  : mask_(m)
144  {}
145 
146  template<typename OptionalArgT>
147  inline value_type eval(int const x, int const y, int const z) const {
148  return mask_(x, y, z);
149  }
150 
151  private:
152  SpatialMask<FieldType> const mask_;
153  };
154  #ifdef __CUDACC__
155  template<typename FieldType>
156  struct NeboMask<GPUWalk, FieldType> {
157  public:
158  FieldType typedef field_type;
159 
160  typename field_type::value_type typedef value_type;
161 
162  NeboMask(int const deviceIndex, SpatialMask<FieldType> const & m)
163  : bitField_(m.mask_values(deviceIndex)),
164  xOffset_(m.window_with_ghost().offset(0) + m.get_ghost_data().get_minus(0)),
165  yOffset_(m.window_with_ghost().offset(1) + m.get_ghost_data().get_minus(1)),
166  zOffset_(m.window_with_ghost().offset(2) + m.get_ghost_data().get_minus(2)),
167  xGlob_(m.window_with_ghost().glob_dim(0)),
168  yGlob_(m.window_with_ghost().glob_dim(1))
169  {}
170 
171  template<typename OptionalArgT>
172  __device__ inline bool eval(int const x, int const y, int const z) const {
173  return deref(x, y, z);
174  }
175 
176  private:
177  __device__ inline int find_position(int const x,
178  int const y,
179  int const z) const {
180  const int newX = xOffset_ + x;
181 
182  const int newY = yOffset_ + y;
183 
184  const int newZ = zOffset_ + z;
185 
186  return newX + xGlob_ * (newY + yGlob_ * newZ);
187  }
188 
189  __device__ inline int find_block(int const position) const {
190  return position / NEBO_INT_BIT;
191  }
192 
193  __device__ inline int find_bit_position(int const position) const {
194  return position % NEBO_INT_BIT;
195  }
196 
197  __device__ inline int deref(int const x, int const y, int const z) const {
198  const int position = find_position(x, y, z);
199 
200  return !(!(*(bitField_ + find_block(position)) & (1 <<
201  find_bit_position(position))));
202  }
203 
204  unsigned int const * bitField_;
205 
206  int const xOffset_;
207 
208  int const yOffset_;
209 
210  int const zOffset_;
211 
212  int const xGlob_;
213 
214  int const yGlob_;
215  }
216  #endif
217  /* __CUDACC__ */;
218  } /* SpatialOps */
219 
220 #endif
221 /* NEBO_MASK_H */
Abstracts a mask.
Definition: SpatialMask.h:70
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...