Reference documentation for deal.II version Git 0994510114 2020-11-30 18:32:56 -0700
\(\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 - 2018 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_COMPILER_CUDA_AWARE
22 
24 
25 namespace LinearAlgebra
26 {
27  namespace CUDAWrappers
28  {
36  inline __device__ float
37  atomicAdd_wrapper(float *address, float val)
38  {
39  return atomicAdd(address, val);
40  }
41 
42 
43 
51  inline __device__ double
52  atomicAdd_wrapper(double *address, double val)
53  {
54  return atomicAdd(address, val);
55  }
56 
57 
58 
64  inline __device__ float
65  atomicMax_wrapper(float *address, float val)
66  {
67  int *address_as_int = reinterpret_cast<int *>(address);
68  int old = *address_as_int, assumed;
69  do
70  {
71  assumed = old;
72  old = atomicCAS(address_as_int,
73  assumed,
74  atomicMax(address_as_int, __float_as_int(val)));
75  }
76  while (assumed != old);
77 
78  return __longlong_as_double(old);
79  }
80 
81 
82 
88  inline __device__ double
89  atomicMax_wrapper(double *address, double val)
90  {
91  unsigned long long int *address_as_ull =
92  reinterpret_cast<unsigned long long int *>(address);
93  unsigned long long int old = *address_as_ull, assumed;
94  do
95  {
96  assumed = old;
97  old = atomicCAS(address_as_ull,
98  assumed,
99  atomicMax(address_as_ull,
100  static_cast<unsigned long long int>(
101  __double_as_longlong(val))));
102  }
103  while (assumed != old);
104 
105  return __longlong_as_double(old);
106  }
107  } // namespace CUDAWrappers
108 } // namespace LinearAlgebra
109 
111 
112 #endif
113 
114 #endif
float atomicMax_wrapper(float *address, float val)
Definition: cuda_atomic.h:65
float atomicAdd_wrapper(float *address, float val)
Definition: cuda_atomic.h:37
#define DEAL_II_NAMESPACE_CLOSE
Definition: config.h:372
#define DEAL_II_NAMESPACE_OPEN
Definition: config.h:371
#define DEAL_II_DEPRECATED
Definition: config.h:152