SpatialOps
NeboReductions.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_REDUCTIONS_H
26  #define NEBO_REDUCTIONS_H
27 
28  #define field_max nebo_max
29 
30  #define field_max_interior nebo_max_interior
31 
32  #define field_min nebo_min
33 
34  #define field_min_interior nebo_min_interior
35 
36  #define field_sum nebo_sum
37 
38  #define field_sum_interior nebo_sum_interior
39 
40  #define field_norm nebo_norm
41 
42  #define field_norm_interior nebo_norm_interior
43 
44  namespace SpatialOps {
45  template<typename ExprType, typename FieldType>
46  inline typename FieldType::value_type nebo_fold(bool useGhost,
47  typename FieldType::
48  value_type (*proc)(typename
49  FieldType::
50  value_type,
51  typename
52  FieldType::
53  value_type),
54  typename FieldType::
55  value_type initialValue,
56  NeboExpression<ExprType,
57  FieldType>
58  const & fexpr) {
59  typename FieldType::value_type typedef value_type;
60 
61  ExprType const initial = fexpr.expr();
62 
63  GhostData const ghosts = (useGhost ? initial.ghosts_without_bc() :
64  GhostData(0));
65 
66  IntVec const extents = initial.extents();
67 
68  IntVec const hasBC = initial.has_bc(PLUS_SIDE);
69 
70  const int xLow = - ghosts.get_minus(0);
71 
72  const int xHigh = extents[0] + ghosts.get_plus(0);
73 
74  const int yLow = - ghosts.get_minus(1);
75 
76  const int yHigh = extents[1] + ghosts.get_plus(1);
77 
78  const int zLow = - ghosts.get_minus(2);
79 
80  const int zHigh = extents[2] + ghosts.get_plus(2);
81 
82  value_type result = initialValue;
83 
84  PoolAutoPtr<int> curEval (Pool<int>::get (CPU_INDEX, 3), CPU_INDEX);
85  double mappedValue;
86 
87  NeboOptionalArg outerIndex(curEval, &mappedValue);
88 
89  typename ExprType::SeqWalkType expr = initial.init(extents,
90  ghosts,
91  hasBC,
92  outerIndex);
93 
94  for(int z = zLow; z < zHigh; z++) {
95  for(int y = yLow; y < yHigh; y++) {
96  for(int x = xLow; x < xHigh; x++) {
97  outerIndex.setOuterIndex(x,y,z);
98 
99  result = proc(result,
100  expr.template eval</*
101  TODO: Wrap Into Compile Time Optional Args
102  */void>(x, y, z));
103  };
104  };
105  };
106 
107  return result;
108  };
109 
110  template<typename AtomicType>
111  inline AtomicType nebo_scalar_max(AtomicType a, AtomicType b) {
112  return ((a < b) ? b : a);
113  };
114 
115  template<typename AtomicType>
116  inline AtomicType nebo_scalar_min(AtomicType a, AtomicType b) {
117  return ((a > b) ? b : a);
118  };
119 
120  template<typename AtomicType>
121  inline AtomicType nebo_scalar_sum(AtomicType a, AtomicType b) {
122  return a + b;
123  };
124 
125  template<typename ExprType, typename FieldType>
126  inline typename FieldType::value_type nebo_max(NeboExpression<ExprType,
127  FieldType>
128  const & fexpr) {
129  return nebo_fold(true,
130  nebo_scalar_max,
131  -(std::numeric_limits<double>().infinity()),
132  fexpr);
133  };
134 
135  template<typename FieldType>
136  inline typename FieldType::value_type nebo_max(FieldType const & field) {
137  NeboConstField<Initial, FieldType> typedef ExprType;
138 
139  return nebo_max(NeboExpression<ExprType, FieldType>(ExprType(field)));
140  };
141 
142  template<typename ExprType, typename FieldType>
143  inline typename FieldType::value_type nebo_max_interior(NeboExpression<ExprType,
144  FieldType>
145  const & fexpr) {
146  return nebo_fold(false,
147  nebo_scalar_max,
148  -(std::numeric_limits<double>().infinity()),
149  fexpr);
150  };
151 
152  template<typename FieldType>
153  inline typename FieldType::value_type nebo_max_interior(FieldType const &
154  field) {
155  NeboConstField<Initial, FieldType> typedef ExprType;
156 
157  return nebo_max_interior(NeboExpression<ExprType, FieldType>(ExprType(field)));
158  };
159 
160  template<typename ExprType, typename FieldType>
161  inline typename FieldType::value_type nebo_min(NeboExpression<ExprType,
162  FieldType>
163  const & fexpr) {
164  return nebo_fold(true,
165  nebo_scalar_min,
166  std::numeric_limits<double>().infinity(),
167  fexpr);
168  };
169 
170  template<typename FieldType>
171  inline typename FieldType::value_type nebo_min(FieldType const & field) {
172  NeboConstField<Initial, FieldType> typedef ExprType;
173 
174  return nebo_min(NeboExpression<ExprType, FieldType>(ExprType(field)));
175  };
176 
177  template<typename ExprType, typename FieldType>
178  inline typename FieldType::value_type nebo_min_interior(NeboExpression<ExprType,
179  FieldType>
180  const & fexpr) {
181  return nebo_fold(false,
182  nebo_scalar_min,
183  std::numeric_limits<double>().infinity(),
184  fexpr);
185  };
186 
187  template<typename FieldType>
188  inline typename FieldType::value_type nebo_min_interior(FieldType const &
189  field) {
190  NeboConstField<Initial, FieldType> typedef ExprType;
191 
192  return nebo_min_interior(NeboExpression<ExprType, FieldType>(ExprType(field)));
193  };
194 
195  template<typename ExprType, typename FieldType>
196  inline typename FieldType::value_type nebo_sum(NeboExpression<ExprType,
197  FieldType>
198  const & fexpr) {
199  return nebo_fold(true, nebo_scalar_sum, 0.0, fexpr);
200  };
201 
202  template<typename FieldType>
203  inline typename FieldType::value_type nebo_sum(FieldType const & field) {
204  NeboConstField<Initial, FieldType> typedef ExprType;
205 
206  return nebo_sum(NeboExpression<ExprType, FieldType>(ExprType(field)));
207  };
208 
209  template<typename ExprType, typename FieldType>
210  inline typename FieldType::value_type nebo_sum_interior(NeboExpression<ExprType,
211  FieldType>
212  const & fexpr) {
213  return nebo_fold(false, nebo_scalar_sum, 0.0, fexpr);
214  };
215 
216  template<typename FieldType>
217  inline typename FieldType::value_type nebo_sum_interior(FieldType const &
218  field) {
219  NeboConstField<Initial, FieldType> typedef ExprType;
220 
221  return nebo_sum_interior(NeboExpression<ExprType, FieldType>(ExprType(field)));
222  };
223 
224  template<typename ExprType, typename FieldType>
225  inline typename FieldType::value_type nebo_norm(NeboExpression<ExprType,
226  FieldType>
227  const & fexpr) {
228  return std::sqrt(nebo_sum(fexpr*fexpr));
229  };
230 
231  template<typename FieldType>
232  inline typename FieldType::value_type nebo_norm(FieldType const & field) {
233  NeboConstField<Initial, FieldType> typedef ExprType;
234 
235  return nebo_norm(NeboExpression<ExprType, FieldType>(ExprType(field)));
236  };
237 
238  template<typename ExprType, typename FieldType>
239  inline typename FieldType::value_type nebo_norm_interior(NeboExpression<ExprType,
240  FieldType>
241  const & fexpr) {
242  return std::sqrt(nebo_sum_interior(fexpr*fexpr));
243  };
244 
245  template<typename FieldType>
246  inline typename FieldType::value_type nebo_norm_interior(FieldType const
247  & field) {
248  NeboConstField<Initial, FieldType> typedef ExprType;
249 
250  return nebo_norm_interior(NeboExpression<ExprType, FieldType>(ExprType(field)));
251  };
252 
253 
254  template<typename CurrentMode, typename Operand, typename Mapper>
256  template<typename Operand, typename Mapper>
257  struct NeboMappedReduction<Initial, Operand, Mapper> {
258  public:
260 
261  #ifdef ENABLE_THREADS
263  #endif
264  /* ENABLE_THREADS */
265 
266  #ifdef __CUDACC__
268  #endif
269  /* __CUDACC__ */
270 
271  NeboMappedReduction(Operand const & operand, Mapper mapper)
272  : operand_(operand), mapper_(mapper)
273  {}
274 
275  inline GhostData ghosts_with_bc(void) const {
276  return operand_.ghosts_with_bc();
277  }
278 
279  inline GhostData ghosts_without_bc(void) const {
280  return operand_.ghosts_without_bc();
281  }
282 
283  inline bool has_extents(void) const {
284  return (operand_.has_extents());
285  }
286 
287  inline IntVec extents(void) const { return operand_.extents(); }
288 
289  inline IntVec has_bc(BCSide const bcSide) const { return operand_.has_bc(bcSide); }
290 
291  inline SeqWalkType init(IntVec const & extents,
292  GhostData const & ghosts,
293  IntVec const & hasBC,
294  NeboOptionalArg & optArg) const {
295 
296  return SeqWalkType(operand_.init(extents,ghosts,hasBC,optArg), mapper_.init(optArg));
297  }
298 
299  #ifdef ENABLE_THREADS
300  inline ResizeType resize(void) const {
301  return ResizeType(operand_.resize(), mapper_.resize());
302  }
303  #endif
304  /* ENABLE_THREADS */
305 
306  #ifdef __CUDACC__
307  inline bool cpu_ready(void) const { return (operand_.cpu_ready()); }
308 
309  inline bool gpu_ready(int const deviceIndex) const {
310  return (operand_.gpu_ready(deviceIndex));
311  }
312 
313  inline GPUWalkType gpu_init(IntVec const & extents,
314  GhostData const & ghosts,
315  IntVec const & hasBC,
316  int const deviceIndex,
317  cudaStream_t const & lhsStream,
318  NeboOptionalArg & optArg) const {
319  return GPUWalkType(operand_.gpu_init(extents,
320  ghosts,
321  hasBC,
322  deviceIndex,
323  lhsStream,
324  optArg), mapper_.gpu_init(deviceIndex, lhsStream, optArg));
325  }
326 
327  inline void stream_wait_event(cudaEvent_t const & event) const {
328  operand_.stream_wait_event(event);
329  }
330 
331  #ifdef NEBO_GPU_TEST
332  inline void gpu_prep(int const deviceIndex) const {
333  operand_.gpu_prep(deviceIndex);
334  }
335  #endif
336  /* NEBO_GPU_TEST */
337  #endif
338  /* __CUDACC__ */
339 
340  private:
341  Operand const operand_;
342  Mapper mapper_;
343  };
344  #ifdef ENABLE_THREADS
345  template<typename Operand, typename Mapper>
346  struct NeboMappedReduction<Resize, Operand, Mapper> {
347  public:
349 
350  NeboMappedReduction(Operand const & operand, Mapper mapper)
351  : operand_(operand), mapper_(mapper)
352  {}
353 
354  inline SeqWalkType init(IntVec const & extents,
355  GhostData const & ghosts,
356  IntVec const & hasBC,
357  NeboOptionalArg & optArg) const {
358 
359  return SeqWalkType(operand_.init(extents, ghosts, hasBC, optArg), mapper_.init(optArg));
360  }
361 
362  private:
363  Operand const operand_;
364  Mapper mapper_;
365  }
366  #endif
367  /* ENABLE_THREADS */;
368  template<typename Operand, typename Mapper>
369  struct NeboMappedReduction<SeqWalk, Operand, Mapper> {
370  public:
371  typename Operand::value_type typedef value_type;
372  typename Mapper::iterator typedef iterator;
373 
374  NeboMappedReduction(Operand const & operand,
375  Mapper mapper)
376  : mapper_(mapper),
377  operand_(operand)
378  { }
379 
380  template<typename OptionalArgT>
381  inline value_type eval(int const x, int const y, int const z) {
382 
383  int it[2];
384 
385  mapper_.eval (x,y,z,it);
386 
387  value_type sum = 0;
388 
389  while (it[1] != -1)
390  {
391  //for summation or generalize?
392  sum += operand_.template eval<OptionalArgT>(it[0],0,0);
393  mapper_.next(x,y,z,it);
394  }
395 
396  return sum;
397  }
398 
399  private:
400  Operand operand_;
401  Mapper mapper_;
402  };
403  #ifdef __CUDACC__
404  template<typename Operand, typename Mapper>
405  struct NeboMappedReduction<GPUWalk, Operand, Mapper> {
406  public:
407  typename Operand::value_type typedef value_type;
408  typename Mapper::iterator typedef iterator;
409 
410  NeboMappedReduction(Operand const & operand, Mapper mapper)
411  : operand_(operand), mapper_(mapper)
412  {}
413 
414  template<typename OptionalArgT>
415  __device__ inline value_type eval(int const x,
416  int const y,
417  int const z) {
418  int it[3];
419 
420  mapper_.eval (x,y,z,it);
421 
422  value_type sum = 0;
423 
424  while (it[1] != -1)
425  {
426  //for summation or generalize?
427  sum += operand_.template eval<OptionalArgT>(it[0],0,0);
428  mapper_.next(x,y,z,it);
429  }
430 
431  return sum;
432 
433  }
434 
435  private:
436  Operand operand_;
437  Mapper mapper_;
438  }
439  #endif
440  /* __CUDACC__ */;
441 
442  /* Field */
443  template<typename FieldType, typename MapperState>
444  inline NeboExpression<NeboMappedReduction<Initial,
445  NeboConstField<Initial,
446  typename NeboFieldCheck<typename
447  FieldType::
448  field_type,
449  FieldType>::
450  Result>,
451  NeboIndexMapper<Initial,
452  MapperState> >,
453  FieldType> nebo_mapped_reduction(FieldType const & arg, const MapperState & mapper) {
454 
456  ReturnType;
457 
458  NeboExpression<ReturnType, FieldType> typedef ReturnTerm;
459 
460  return ReturnTerm(ReturnType(NeboConstField<Initial, FieldType>(arg), NeboIndexMapper<Initial, MapperState> (mapper)));
461  }
462 
463  /* SubExpr */
464  template<typename SubExpr, typename MapperState, typename FieldType>
465  inline NeboExpression<NeboMappedReduction<Initial, SubExpr,
466  NeboIndexMapper<Initial,
467  MapperState> >,
468  FieldType> nebo_mapped_reduction(NeboExpression<SubExpr,
469  FieldType>
470  const & arg, const MapperState & mapper) {
472  ReturnType;
473 
474  NeboExpression<ReturnType, FieldType> typedef ReturnTerm;
475 
476  return ReturnTerm(ReturnType(arg.expr(), mapper));
477  }
478 
479  } /* SpatialOps */
480 
481 #endif
482 /* NEBO_REDUCTIONS_H */
483 
static T * get(const short int deviceLocation, const size_t n)
Obtain a pointer to a block of memory with the requested size. If the requested size is zero...
Definition: MemoryPool.cpp:116
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...