Reference documentation for deal.II version 9.5.0
\(\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\}}\)
Loading...
Searching...
No Matches
cuda_atomic.h
Go to the documentation of this file.
1// ---------------------------------------------------------------------
2//
3// Copyright (C) 2016 - 2022 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
25namespace 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:472
#define DEAL_II_NAMESPACE_CLOSE
Definition config.h:473
float atomicMax_wrapper(float *address, float val)
Definition cuda_atomic.h:35