Loading [MathJax]/extensions/TeX/AMSsymbols.js
 Reference documentation for deal.II version 9.4.1
\(\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\}}\)
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Modules Pages
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_COMPILER_CUDA_AWARE
22
24
25namespace 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
#define DEAL_II_DEPRECATED
Definition: config.h:164
#define DEAL_II_NAMESPACE_OPEN
Definition: config.h:442
#define DEAL_II_NAMESPACE_CLOSE
Definition: config.h:443
float atomicMax_wrapper(float *address, float val)
Definition: cuda_atomic.h:65
float atomicAdd_wrapper(float *address, float val)
Definition: cuda_atomic.h:37