Sacado Package Browser (Single Doxygen Collection)  Version of the Day
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Sacado_Fad_Exp_Atomic.hpp
Go to the documentation of this file.
1 // @HEADER
2 // *****************************************************************************
3 // Sacado Package
4 //
5 // Copyright 2006 NTESS and the Sacado contributors.
6 // SPDX-License-Identifier: LGPL-2.1-or-later
7 // *****************************************************************************
8 // @HEADER
9 
10 #ifndef SACADO_FAD_EXP_ATOMIC_HPP
11 #define SACADO_FAD_EXP_ATOMIC_HPP
12 
13 #include "Sacado_ConfigDefs.h"
14 #if defined(HAVE_SACADO_KOKKOS)
15 
16 #include "Sacado_Fad_Exp_ViewFad.hpp"
17 #include "Kokkos_Atomic.hpp"
18 #include "impl/Kokkos_Error.hpp"
19 
20 namespace Sacado {
21 
22  namespace Fad {
23  namespace Exp {
24 
25  // Overload of Kokkos::atomic_add for ViewFad types.
26  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
27  SACADO_INLINE_FUNCTION
28  void atomic_add(ViewFadPtr<ValT,sl,ss,U> dst, const Expr<T>& xx) {
29  using Kokkos::atomic_add;
30 
31  const typename Expr<T>::derived_type& x = xx.derived();
32 
33  const int xsz = x.size();
34  const int sz = dst->size();
35 
36  // We currently cannot handle resizing since that would need to be
37  // done atomically.
38  if (xsz > sz)
39  Kokkos::abort(
40  "Sacado error: Fad resize within atomic_add() not supported!");
41 
42  if (xsz != sz && sz > 0 && xsz > 0)
43  Kokkos::abort(
44  "Sacado error: Fad assignment of incompatiable sizes!");
45 
46 
47  if (sz > 0 && xsz > 0) {
48  SACADO_FAD_DERIV_LOOP(i,sz)
49  atomic_add(&(dst->fastAccessDx(i)), x.fastAccessDx(i));
50  }
51  SACADO_FAD_THREAD_SINGLE
52  atomic_add(&(dst->val()), x.val());
53  }
54 
55  namespace Impl {
56  // Our implementation of Kokkos::atomic_oper_fetch() and
57  // Kokkos::atomic_fetch_oper() for Sacado types on host
58  template <typename Oper, typename DestPtrT, typename ValT, typename T>
59  typename Sacado::BaseExprType< Expr<T> >::type
60  atomic_oper_fetch_host(const Oper& op, DestPtrT dest, ValT* dest_val,
61  const Expr<T>& x)
62  {
63  typedef typename Sacado::BaseExprType< Expr<T> >::type return_type;
64  const typename Expr<T>::derived_type& val = x.derived();
65 
66 #ifdef KOKKOS_INTERNAL_NOT_PARALLEL
67  auto scope = desul::MemoryScopeCaller();
68 #else
69  auto scope = desul::MemoryScopeDevice();
70 #endif
71 
72  while (!desul::Impl::lock_address((void*)dest_val, scope))
73  ;
74  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
75  return_type return_val = op.apply(*dest, val);
76  *dest = return_val;
77  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
78  desul::Impl::unlock_address((void*)dest_val, scope);
79  return return_val;
80  }
81 
82  template <typename Oper, typename DestPtrT, typename ValT, typename T>
83  typename Sacado::BaseExprType< Expr<T> >::type
84  atomic_fetch_oper_host(const Oper& op, DestPtrT dest, ValT* dest_val,
85  const Expr<T>& x)
86  {
87  typedef typename Sacado::BaseExprType< Expr<T> >::type return_type;
88  const typename Expr<T>::derived_type& val = x.derived();
89 
90 #ifdef KOKKOS_INTERNAL_NOT_PARALLEL
91  auto scope = desul::MemoryScopeCaller();
92 #else
93  auto scope = desul::MemoryScopeDevice();
94 #endif
95 
96  while (!desul::Impl::lock_address((void*)dest_val, scope))
97  ;
98  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
99  return_type return_val = *dest;
100  *dest = op.apply(return_val, val);
101  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
102  desul::Impl::unlock_address((void*)dest_val, scope);
103  return return_val;
104  }
105 
106  // Helper function to decide if we are using team-based parallelism
107 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
108  __device__
109  inline bool atomics_use_team() {
110 #if defined(SACADO_VIEW_CUDA_HIERARCHICAL) || defined(SACADO_VIEW_CUDA_HIERARCHICAL_DFAD)
111  // It is not allowed to define SACADO_VIEW_CUDA_HIERARCHICAL or
112  // SACADO_VIEW_CUDA_HIERARCHICAL_DFAD and use Sacado inside a team-based
113  // kernel without Sacado hierarchical parallelism. So use the
114  // team-based version only if blockDim.x > 1 (i.e., a team policy)
115  return (blockDim.x > 1);
116 #else
117  return false;
118 #endif
119  }
120 #endif
121 
122 #if defined(KOKKOS_ENABLE_CUDA)
123 
124  // Our implementation of Kokkos::atomic_oper_fetch() and
125  // Kokkos::atomic_fetch_oper() for Sacado types on device
126  template <typename Oper, typename DestPtrT, typename ValT, typename T>
127  __device__
128  typename Sacado::BaseExprType< Expr<T> >::type
129  atomic_oper_fetch_device(const Oper& op, DestPtrT dest, ValT* dest_val,
130  const Expr<T>& x)
131  {
132  typedef typename Sacado::BaseExprType< Expr<T> >::type return_type;
133  const typename Expr<T>::derived_type& val = x.derived();
134 
135  auto scope = desul::MemoryScopeDevice();
136 
137  if (atomics_use_team()) {
138  int go = 1;
139  while (go) {
140  if (threadIdx.x == 0)
141  go = !desul::Impl::lock_address_cuda((void*)dest_val, scope);
142  go = Kokkos::shfl(go, 0, blockDim.x);
143  }
144  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
145  return_type return_val = op.apply(*dest, val);
146  *dest = return_val;
147  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
148  if (threadIdx.x == 0)
149  desul::Impl::unlock_address_cuda((void*)dest_val, scope);
150  return return_val;
151  }
152  else {
153  return_type return_val;
154  // This is a way to avoid dead lock in a warp
155  int done = 0;
156  unsigned int mask = __activemask() ;
157  unsigned int active = __ballot_sync(mask, 1);
158  unsigned int done_active = 0;
159  while (active != done_active) {
160  if (!done) {
161  if (desul::Impl::lock_address_cuda((void*)dest_val, scope)) {
162  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
163  return_val = op.apply(*dest, val);
164  *dest = return_val;
165  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
166  desul::Impl::unlock_address_cuda((void*)dest_val, scope);
167  done = 1;
168  }
169  }
170  done_active = __ballot_sync(mask, done);
171  }
172  return return_val;
173  }
174  }
175 
176  template <typename Oper, typename DestPtrT, typename ValT, typename T>
177  __device__
178  typename Sacado::BaseExprType< Expr<T> >::type
179  atomic_fetch_oper_device(const Oper& op, DestPtrT dest, ValT* dest_val,
180  const Expr<T>& x)
181  {
182  typedef typename Sacado::BaseExprType< Expr<T> >::type return_type;
183  const typename Expr<T>::derived_type& val = x.derived();
184 
185  auto scope = desul::MemoryScopeDevice();
186 
187  if (atomics_use_team()) {
188  int go = 1;
189  while (go) {
190  if (threadIdx.x == 0)
191  go = !desul::Impl::lock_address_cuda((void*)dest_val, scope);
192  go = Kokkos::shfl(go, 0, blockDim.x);
193  }
194  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
195  return_type return_val = *dest;
196  *dest = op.apply(return_val, val);
197  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
198  if (threadIdx.x == 0)
199  desul::Impl::unlock_address_cuda((void*)dest_val, scope);
200  return return_val;
201  }
202  else {
203  return_type return_val;
204  // This is a way to (hopefully) avoid dead lock in a warp
205  int done = 0;
206  unsigned int mask = __activemask() ;
207  unsigned int active = __ballot_sync(mask, 1);
208  unsigned int done_active = 0;
209  while (active != done_active) {
210  if (!done) {
211  if (desul::Impl::lock_address_cuda((void*)dest_val, scope)) {
212  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
213  return_val = *dest;
214  *dest = op.apply(return_val, val);
215  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
216  desul::Impl::unlock_address_cuda((void*)dest_val, scope);
217  done = 1;
218  }
219  }
220  done_active = __ballot_sync(mask, done);
221  }
222  return return_val;
223  }
224  }
225 
226 #elif defined(KOKKOS_ENABLE_HIP)
227 
228  // Our implementation of Kokkos::atomic_oper_fetch() and
229  // Kokkos::atomic_fetch_oper() for Sacado types on device
230  template <typename Oper, typename DestPtrT, typename ValT, typename T>
231  __device__
232  typename Sacado::BaseExprType< Expr<T> >::type
233  atomic_oper_fetch_device(const Oper& op, DestPtrT dest, ValT* dest_val,
234  const Expr<T>& x)
235  {
236  typedef typename Sacado::BaseExprType< Expr<T> >::type return_type;
237  const typename Expr<T>::derived_type& val = x.derived();
238 
239  auto scope = desul::MemoryScopeDevice();
240 
241  if (atomics_use_team()) {
242  int go = 1;
243  while (go) {
244  if (threadIdx.x == 0)
245  go = !desul::Impl::lock_address_hip((void*)dest_val, scope);
246  go = Kokkos::shfl(go, 0, blockDim.x);
247  }
248  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
249  return_type return_val = op.apply(*dest, val);
250  *dest = return_val;
251  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
252  if (threadIdx.x == 0)
253  desul::Impl::unlock_address_hip((void*)dest_val, scope);
254  return return_val;
255  }
256  else {
257  return_type return_val;
258  int done = 0;
259  unsigned int active = __ballot(1);
260  unsigned int done_active = 0;
261  while (active != done_active) {
262  if (!done) {
263  if (desul::Impl::lock_address_hip((void*)dest_val, scope)) {
264  return_val = op.apply(*dest, val);
265  *dest = return_val;
266  desul::Impl::unlock_address_hip((void*)dest_val, scope);
267  done = 1;
268  }
269  }
270  done_active = __ballot(done);
271  }
272  return return_val;
273  }
274  }
275 
276  template <typename Oper, typename DestPtrT, typename ValT, typename T>
277  __device__
278  typename Sacado::BaseExprType< Expr<T> >::type
279  atomic_fetch_oper_device(const Oper& op, DestPtrT dest, ValT* dest_val,
280  const Expr<T>& x)
281  {
282  typedef typename Sacado::BaseExprType< Expr<T> >::type return_type;
283  const typename Expr<T>::derived_type& val = x.derived();
284 
285  auto scope = desul::MemoryScopeDevice();
286 
287  if (atomics_use_team()) {
288  int go = 1;
289  while (go) {
290  if (threadIdx.x == 0)
291  go = !desul::Impl::lock_address_hip((void*)dest_val, scope);
292  go = Kokkos::shfl(go, 0, blockDim.x);
293  }
294  desul::atomic_thread_fence(desul::MemoryOrderAcquire(), scope);
295  return_type return_val = *dest;
296  *dest = op.apply(return_val, val);
297  desul::atomic_thread_fence(desul::MemoryOrderRelease(), scope);
298  if (threadIdx.x == 0)
299  desul::Impl::unlock_address_hip((void*)dest_val, scope);
300  return return_val;
301  }
302  else {
303  return_type return_val;
304  int done = 0;
305  unsigned int active = __ballot(1);
306  unsigned int done_active = 0;
307  while (active != done_active) {
308  if (!done) {
309  if (desul::Impl::lock_address_hip((void*)dest_val, scope)) {
310  return_val = *dest;
311  *dest = op.apply(return_val, val);
312  desul::Impl::unlock_address_hip((void*)dest_val, scope);
313  done = 1;
314  }
315  }
316  done_active = __ballot(done);
317  }
318  return return_val;
319  }
320  }
321 
322 #elif defined(KOKKOS_ENABLE_SYCL)
323 
324  // Our implementation of Kokkos::atomic_oper_fetch() and
325  // Kokkos::atomic_fetch_oper() for Sacado types on device
326  template <typename Oper, typename DestPtrT, typename ValT, typename T>
327  typename Sacado::BaseExprType< Expr<T> >::type
328  atomic_oper_fetch_device(const Oper& op, DestPtrT dest, ValT* dest_val,
329  const Expr<T>& x)
330  {
331  Kokkos::abort("Not implemented!");
332  return {};
333  }
334 
335  template <typename Oper, typename DestPtrT, typename ValT, typename T>
336  typename Sacado::BaseExprType< Expr<T> >::type
337  atomic_fetch_oper_device(const Oper& op, DestPtrT dest, ValT* dest_val,
338  const Expr<T>& x)
339  {
340  Kokkos::abort("Not implemented!");
341  return {};
342  }
343 #endif
344 
345  // Overloads of Kokkos::atomic_oper_fetch/Kokkos::atomic_fetch_oper
346  // for Sacado types
347  template <typename Oper, typename S>
348  SACADO_INLINE_FUNCTION GeneralFad<S>
349  atomic_oper_fetch(const Oper& op, GeneralFad<S>* dest,
350  const GeneralFad<S>& val)
351  {
352  KOKKOS_IF_ON_HOST(return Impl::atomic_oper_fetch_host(op, dest, &(dest->val()), val);)
353  KOKKOS_IF_ON_DEVICE(return Impl::atomic_oper_fetch_device(op, dest, &(dest->val()), val);)
354  }
355  template <typename Oper, typename ValT, unsigned sl, unsigned ss,
356  typename U, typename T>
357  SACADO_INLINE_FUNCTION U
358  atomic_oper_fetch(const Oper& op, ViewFadPtr<ValT,sl,ss,U> dest,
359  const Expr<T>& val)
360  {
361  KOKKOS_IF_ON_HOST(return Impl::atomic_oper_fetch_host(op, dest, &dest.val(), val);)
362  KOKKOS_IF_ON_DEVICE(return Impl::atomic_oper_fetch_device(op, dest, &dest.val(), val);)
363  }
364 
365  template <typename Oper, typename S>
366  SACADO_INLINE_FUNCTION GeneralFad<S>
367  atomic_fetch_oper(const Oper& op, GeneralFad<S>* dest,
368  const GeneralFad<S>& val)
369  {
370  KOKKOS_IF_ON_HOST(return Impl::atomic_fetch_oper_host(op, dest, &(dest->val()), val);)
371  KOKKOS_IF_ON_DEVICE(return Impl::atomic_fetch_oper_device(op, dest, &(dest->val()), val);)
372  }
373  template <typename Oper, typename ValT, unsigned sl, unsigned ss,
374  typename U, typename T>
375  SACADO_INLINE_FUNCTION U
376  atomic_fetch_oper(const Oper& op, ViewFadPtr<ValT,sl,ss,U> dest,
377  const Expr<T>& val)
378  {
379  KOKKOS_IF_ON_HOST(return Impl::atomic_fetch_oper_host(op, dest, &dest.val(), val);)
380  KOKKOS_IF_ON_DEVICE(return Impl::atomic_fetch_oper_device(op, dest, &dest.val(), val);)
381  }
382 
383  // Our definition of the various Oper classes to be more type-flexible
384  struct MaxOper {
385  template <class Scalar1, class Scalar2>
386  KOKKOS_FORCEINLINE_FUNCTION
387  static auto apply(const Scalar1& val1, const Scalar2& val2)
388  -> decltype(max(val1,val2))
389  {
390  return max(val1,val2);
391  }
392  };
393  struct MinOper {
394  template <class Scalar1, class Scalar2>
395  KOKKOS_FORCEINLINE_FUNCTION
396  static auto apply(const Scalar1& val1, const Scalar2& val2)
397  -> decltype(min(val1,val2))
398  {
399  return min(val1,val2);
400  }
401  };
402  struct AddOper {
403  template <class Scalar1, class Scalar2>
404  KOKKOS_FORCEINLINE_FUNCTION
405  static auto apply(const Scalar1& val1, const Scalar2& val2)
406  -> decltype(val1+val2)
407  {
408  return val1 + val2;
409  }
410  };
411  struct SubOper {
412  template <class Scalar1, class Scalar2>
413  KOKKOS_FORCEINLINE_FUNCTION
414  static auto apply(const Scalar1& val1, const Scalar2& val2)
415  -> decltype(val1-val2)
416  {
417  return val1 - val2;
418  }
419  };
420  struct MulOper {
421  template <class Scalar1, class Scalar2>
422  KOKKOS_FORCEINLINE_FUNCTION
423  static auto apply(const Scalar1& val1, const Scalar2& val2)
424  -> decltype(val1*val2)
425  {
426  return val1 * val2;
427  }
428  };
429  struct DivOper {
430  template <class Scalar1, class Scalar2>
431  KOKKOS_FORCEINLINE_FUNCTION
432  static auto apply(const Scalar1& val1, const Scalar2& val2)
433  -> decltype(val1/val2)
434  {
435  return val1 / val2;
436  }
437  };
438 
439  } // Impl
440 
441  // Overload of Kokkos::atomic_*_fetch() and Kokkos::atomic_fetch_*()
442  // for Sacado types
443  template <typename S>
444  SACADO_INLINE_FUNCTION GeneralFad<S>
445  atomic_max_fetch(GeneralFad<S>* dest, const GeneralFad<S>& val) {
446  return Impl::atomic_oper_fetch(Impl::MaxOper(), dest, val);
447  }
448  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
449  SACADO_INLINE_FUNCTION U
450  atomic_max_fetch(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
451  return Impl::atomic_oper_fetch(Impl::MaxOper(), dest, val);
452  }
453  template <typename S>
454  SACADO_INLINE_FUNCTION GeneralFad<S>
455  atomic_min_fetch(GeneralFad<S>* dest, const GeneralFad<S>& val) {
456  return Impl::atomic_oper_fetch(Impl::MinOper(), dest, val);
457  }
458  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
459  SACADO_INLINE_FUNCTION U
460  atomic_min_fetch(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
461  return Impl::atomic_oper_fetch(Impl::MinOper(), dest, val);
462  }
463  template <typename S>
464  SACADO_INLINE_FUNCTION GeneralFad<S>
465  atomic_add_fetch(GeneralFad<S>* dest, const GeneralFad<S>& val) {
466  return Impl::atomic_oper_fetch(Impl::AddOper(), dest, val);
467  }
468  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
469  SACADO_INLINE_FUNCTION U
470  atomic_add_fetch(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
471  return Impl::atomic_oper_fetch(Impl::AddOper(), dest, val);
472  }
473  template <typename S>
474  SACADO_INLINE_FUNCTION GeneralFad<S>
475  atomic_sub_fetch(GeneralFad<S>* dest, const GeneralFad<S>& val) {
476  return Impl::atomic_oper_fetch(Impl::SubOper(), dest, val);
477  }
478  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
479  SACADO_INLINE_FUNCTION U
480  atomic_sub_fetch(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
481  return Impl::atomic_oper_fetch(Impl::SubOper(), dest, val);
482  }
483  template <typename S>
484  SACADO_INLINE_FUNCTION GeneralFad<S>
485  atomic_mul_fetch(GeneralFad<S>* dest, const GeneralFad<S>& val) {
486  return atomic_oper_fetch(Impl::MulOper(), dest, val);
487  }
488  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
489  SACADO_INLINE_FUNCTION U
490  atomic_mul_fetch(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
491  return Impl::atomic_oper_fetch(Impl::MulOper(), dest, val);
492  }
493  template <typename S>
494  SACADO_INLINE_FUNCTION GeneralFad<S>
495  atomic_div_fetch(GeneralFad<S>* dest, const GeneralFad<S>& val) {
496  return Impl::atomic_oper_fetch(Impl::DivOper(), dest, val);
497  }
498  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
499  SACADO_INLINE_FUNCTION U
500  atomic_div_fetch(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
501  return Impl::atomic_oper_fetch(Impl::DivOper(), dest, val);
502  }
503 
504  template <typename S>
505  SACADO_INLINE_FUNCTION GeneralFad<S>
506  atomic_fetch_max(GeneralFad<S>* dest, const GeneralFad<S>& val) {
507  return Impl::atomic_fetch_oper(Impl::MaxOper(), dest, val);
508  }
509  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
510  SACADO_INLINE_FUNCTION U
511  atomic_fetch_max(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
512  return Impl::atomic_fetch_oper(Impl::MaxOper(), dest, val);
513  }
514  template <typename S>
515  SACADO_INLINE_FUNCTION GeneralFad<S>
516  atomic_fetch_min(GeneralFad<S>* dest, const GeneralFad<S>& val) {
517  return Impl::atomic_fetch_oper(Impl::MinOper(), dest, val);
518  }
519  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
520  SACADO_INLINE_FUNCTION U
521  atomic_fetch_min(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
522  return Impl::atomic_fetch_oper(Impl::MinOper(), dest, val);
523  }
524  template <typename S>
525  SACADO_INLINE_FUNCTION GeneralFad<S>
526  atomic_fetch_add(GeneralFad<S>* dest, const GeneralFad<S>& val) {
527  return Impl::atomic_fetch_oper(Impl::AddOper(), dest, val);
528  }
529  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
530  SACADO_INLINE_FUNCTION U
531  atomic_fetch_add(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
532  return Impl::atomic_fetch_oper(Impl::AddOper(), dest, val);
533  }
534  template <typename S>
535  SACADO_INLINE_FUNCTION GeneralFad<S>
536  atomic_fetch_sub(GeneralFad<S>* dest, const GeneralFad<S>& val) {
537  return Impl::atomic_fetch_oper(Impl::SubOper(), dest, val);
538  }
539  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
540  SACADO_INLINE_FUNCTION U
541  atomic_fetch_sub(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
542  return Impl::atomic_fetch_oper(Impl::SubOper(), dest, val);
543  }
544  template <typename S>
545  SACADO_INLINE_FUNCTION GeneralFad<S>
546  atomic_fetch_mul(GeneralFad<S>* dest, const GeneralFad<S>& val) {
547  return Impl::atomic_fetch_oper(Impl::MulOper(), dest, val);
548  }
549  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
550  SACADO_INLINE_FUNCTION U
551  atomic_fetch_mul(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
552  return Impl::atomic_fetch_oper(Impl::MulOper(), dest, val);
553  }
554  template <typename S>
555  SACADO_INLINE_FUNCTION GeneralFad<S>
556  atomic_fetch_div(GeneralFad<S>* dest, const GeneralFad<S>& val) {
557  return Impl::atomic_fetch_oper(Impl::DivOper(), dest, val);
558  }
559  template <typename ValT, unsigned sl, unsigned ss, typename U, typename T>
560  SACADO_INLINE_FUNCTION U
561  atomic_fetch_div(ViewFadPtr<ValT,sl,ss,U> dest, const Expr<T>& val) {
562  return Impl::atomic_fetch_oper(Impl::DivOper(), dest, val);
563  }
564 
565  } // namespace Exp
566  } // namespace Fad
567 
568 } // namespace Sacado
569 
570 #endif // HAVE_SACADO_KOKKOS
571 #endif // SACADO_FAD_EXP_VIEWFAD_HPP
SACADO_FAD_THREAD_SINGLE
#define SACADO_FAD_THREAD_SINGLE
Definition: Sacado_Fad_Exp_ExprAssign.hpp:30
i
int i
Definition: gmock-matchers_test.cc:718
val
expr val()
T
#define T
Definition: Sacado_rad.hpp:553
Sacado::Fad::min
SimpleFad< ValueT > min(const SimpleFad< ValueT > &a, const SimpleFad< ValueT > &b)
Definition: Sacado_Fad_SimpleFadOps.hpp:479
SACADO_FAD_DERIV_LOOP
#define SACADO_FAD_DERIV_LOOP(I, SZ)
Definition: Sacado_Fad_Exp_ExprAssign.hpp:22
Sacado::BaseExprType
Get the base Fad type from a view/expression.
Definition: Sacado_Traits.hpp:70
fastAccessDx
expr expr expr fastAccessDx(i)) FAD_UNARYOP_MACRO(exp
Definition: Sacado_Fad_Exp_Ops.hpp:267
x
int x
Definition: gmock-matchers_test.cc:3775
Sacado::Fad::max
SimpleFad< ValueT > max(const SimpleFad< ValueT > &a, const SimpleFad< ValueT > &b)
Definition: Sacado_Fad_SimpleFadOps.hpp:415
Sacado_ConfigDefs.h
SACADO_INLINE_FUNCTION
#define SACADO_INLINE_FUNCTION
Definition: Sacado_ConfigDefs.h:88
Sacado_Fad_Exp_ViewFad.hpp
Generated on Tue Nov 12 2024 09:21:50 for Sacado Package Browser (Single Doxygen Collection) by   doxygen 1.8.5