Reference documentation for deal.II version GIT 5983d193e2 2023-05-27 16:50:02+00:00
\(\newcommand{\dealvcentcolon}{\mathrel{\mathop{:}}}\) \(\newcommand{\dealcoloneq}{\dealvcentcolon\mathrel{\mkern-1.2mu}=}\) \(\newcommand{\jump}[1]{\left[\!\left[ #1 \right]\!\right]}\) \(\newcommand{\average}[1]{\left\{\!\left\{ #1 \right\}\!\right\}}\)
cuda_atomic.h
Go to the documentation of this file.
1 // ---------------------------------------------------------------------
2 //
3 // Copyright (C) 2016 - 2020 by the deal.II authors
4 //
5 // This file is part of the deal.II library.
6 //
7 // The deal.II library is free software; you can use it, redistribute
8 // it, and/or modify it under the terms of the GNU Lesser General
9 // Public License as published by the Free Software Foundation; either
10 // version 2.1 of the License, or (at your option) any later version.
11 // The full text of the license can be found in the file LICENSE.md at
12 // the top level directory of deal.II.
13 //
14 // ---------------------------------------------------------------------
15 
16 #ifndef dealii_cuda_atomic_h
17 #define dealii_cuda_atomic_h
18 
19 #include <deal.II/base/config.h>
20 
21 #ifdef DEAL_II_WITH_CUDA
22 
24 
25 namespace LinearAlgebra
26 {
27  namespace CUDAWrappers
28  {
34  inline __device__ float
35  atomicMax_wrapper(float *address, float val)
36  {
37  int *address_as_int = reinterpret_cast<int *>(address);
38  int old = *address_as_int, assumed;
39  do
40  {
41  assumed = old;
42  old = atomicCAS(address_as_int,
43  assumed,
44  atomicMax(address_as_int, __float_as_int(val)));
45  }
46  while (assumed != old);
47 
48  return __longlong_as_double(old);
49  }
50 
51 
52 
58  inline __device__ double
59  atomicMax_wrapper(double *address, double val)
60  {
61  unsigned long long int *address_as_ull =
62  reinterpret_cast<unsigned long long int *>(address);
63  unsigned long long int old = *address_as_ull, assumed;
64  do
65  {
66  assumed = old;
67  old = atomicCAS(address_as_ull,
68  assumed,
69  atomicMax(address_as_ull,
70  static_cast<unsigned long long int>(
71  __double_as_longlong(val))));
72  }
73  while (assumed != old);
74 
75  return __longlong_as_double(old);
76  }
77  } // namespace CUDAWrappers
78 } // namespace LinearAlgebra
79 
81 
82 #endif
83 
84 #endif
#define DEAL_II_NAMESPACE_OPEN
Definition: config.h:474
#define DEAL_II_NAMESPACE_CLOSE
Definition: config.h:475
float atomicMax_wrapper(float *address, float val)
Definition: cuda_atomic.h:35