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,
73  NeboOptionalArg & optArg) const {
74  return SeqWalkType(mask_);
75  }
76 
77  #ifdef ENABLE_THREADS
78  inline ResizeType resize(void) const { return ResizeType(mask_); }
79  #endif
80  /* ENABLE_THREADS */
81 
82  #ifdef __CUDACC__
83  inline bool cpu_ready(void) const {
84  return mask_.find_consumer(CPU_INDEX);
85  }
86 
87  inline bool gpu_ready(int const deviceIndex) const {
88  return mask_.find_consumer(deviceIndex);
89  }
90 
91  inline GPUWalkType gpu_init(IntVec const & extents,
92  GhostData const & ghosts,
93  IntVec const & hasBC,
94  int const deviceIndex,
95  cudaStream_t const & lhsStream,
96  NeboOptionalArg & optArg) const {
97  return GPUWalkType(deviceIndex, mask_);
98  }
99 
100  inline void stream_wait_event(cudaEvent_t const & event) const {}
101 
102  #ifdef NEBO_GPU_TEST
103  inline void gpu_prep(int const deviceIndex) const {
104  const_cast<SpatialMask<FieldType> *>(&mask_)->add_consumer(deviceIndex);
105  }
106  #endif
107  /* NEBO_GPU_TEST */
108  #endif
109  /* __CUDACC__ */
110 
111  private:
112  SpatialMask<FieldType> const mask_;
113  };
114  #ifdef ENABLE_THREADS
115  template<typename FieldType>
116  struct NeboMask<Resize, FieldType> {
117  public:
118  FieldType typedef field_type;
119 
120  NeboMask<SeqWalk, FieldType> typedef SeqWalkType;
121 
123  : mask_(m)
124  {}
125 
126  inline SeqWalkType init(IntVec const & extents,
127  GhostData const & ghosts,
128  IntVec const & hasBC,
129  NeboOptionalArg & optArg) const {
130  return SeqWalkType(mask_);
131  }
132 
133  private:
134  SpatialMask<FieldType> const mask_;
135  }
136  #endif
137  /* ENABLE_THREADS */;
138  template<typename FieldType>
139  struct NeboMask<SeqWalk, FieldType> {
140  public:
141  FieldType typedef field_type;
142 
143  typename field_type::value_type typedef value_type;
144 
146  : mask_(m)
147  {}
148 
149  template<typename OptionalArgT>
150  inline value_type eval(int const x, int const y, int const z) const {
151  return mask_(x, y, z);
152  }
153 
154  private:
155  SpatialMask<FieldType> const mask_;
156  };
157  #ifdef __CUDACC__
158  template<typename FieldType>
159  struct NeboMask<GPUWalk, FieldType> {
160  public:
161  FieldType typedef field_type;
162 
163  typename field_type::value_type typedef value_type;
164 
165  NeboMask(int const deviceIndex, SpatialMask<FieldType> const & m)
166  : bitField_(m.mask_values(deviceIndex)),
167  xOffset_(m.window_with_ghost().offset(0) + m.get_ghost_data().get_minus(0)),
168  yOffset_(m.window_with_ghost().offset(1) + m.get_ghost_data().get_minus(1)),
169  zOffset_(m.window_with_ghost().offset(2) + m.get_ghost_data().get_minus(2)),
170  xGlob_(m.window_with_ghost().glob_dim(0)),
171  yGlob_(m.window_with_ghost().glob_dim(1))
172  {}
173 
174  template<typename OptionalArgT>
175  __device__ inline bool eval(int const x, int const y, int const z) const {
176  return deref(x, y, z);
177  }
178 
179  private:
180  __device__ inline int find_position(int const x,
181  int const y,
182  int const z) const {
183  const int newX = xOffset_ + x;
184 
185  const int newY = yOffset_ + y;
186 
187  const int newZ = zOffset_ + z;
188 
189  return newX + xGlob_ * (newY + yGlob_ * newZ);
190  }
191 
192  __device__ inline int find_block(int const position) const {
193  return position / NEBO_INT_BIT;
194  }
195 
196  __device__ inline int find_bit_position(int const position) const {
197  return position % NEBO_INT_BIT;
198  }
199 
200  __device__ inline int deref(int const x, int const y, int const z) const {
201  const int position = find_position(x, y, z);
202 
203  return !(!(*(bitField_ + find_block(position)) & (1 <<
204  find_bit_position(position))));
205  }
206 
207  unsigned int const * bitField_;
208 
209  int const xOffset_;
210 
211  int const yOffset_;
212 
213  int const zOffset_;
214 
215  int const xGlob_;
216 
217  int const yGlob_;
218  }
219  #endif
220  /* __CUDACC__ */;
221  } /* SpatialOps */
222 
223 #endif
224 /* 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
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...