Reference documentation for deal.II version GIT relicensing-437-g81ec864850 2024-04-19 07:30: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\}}\)
Loading...
Searching...
No Matches
cuda_atomic.h
Go to the documentation of this file.
1// ------------------------------------------------------------------------
2//
3// SPDX-License-Identifier: LGPL-2.1-or-later
4// Copyright (C) 2016 - 2022 by the deal.II authors
5//
6// This file is part of the deal.II library.
7//
8// Part of the source code is dual licensed under Apache-2.0 WITH
9// LLVM-exception OR LGPL-2.1-or-later. Detailed license information
10// governing the source code and code contributions can be found in
11// LICENSE.md and CONTRIBUTING.md at the top level directory of deal.II.
12//
13// ------------------------------------------------------------------------
14
15#ifndef dealii_cuda_atomic_h
16#define dealii_cuda_atomic_h
17
18#include <deal.II/base/config.h>
19
20#ifdef DEAL_II_WITH_CUDA
21
23
24namespace LinearAlgebra
25{
26 namespace CUDAWrappers
27 {
33 inline __device__ float
34 atomicMax_wrapper(float *address, float val)
35 {
36 int *address_as_int = reinterpret_cast<int *>(address);
37 int old = *address_as_int, assumed;
38 do
39 {
40 assumed = old;
41 old = atomicCAS(address_as_int,
42 assumed,
43 atomicMax(address_as_int, __float_as_int(val)));
44 }
45 while (assumed != old);
46
47 return __longlong_as_double(old);
48 }
49
50
51
57 inline __device__ double
58 atomicMax_wrapper(double *address, double val)
59 {
60 unsigned long long int *address_as_ull =
61 reinterpret_cast<unsigned long long int *>(address);
62 unsigned long long int old = *address_as_ull, assumed;
63 do
64 {
65 assumed = old;
66 old = atomicCAS(address_as_ull,
67 assumed,
68 atomicMax(address_as_ull,
69 static_cast<unsigned long long int>(
70 __double_as_longlong(val))));
71 }
72 while (assumed != old);
73
74 return __longlong_as_double(old);
75 }
76 } // namespace CUDAWrappers
77} // namespace LinearAlgebra
78
80
81#endif
82
83#endif
#define DEAL_II_NAMESPACE_OPEN
Definition config.h:502
#define DEAL_II_NAMESPACE_CLOSE
Definition config.h:503
float atomicMax_wrapper(float *address, float val)
Definition cuda_atomic.h:34