File: fetch_max.md

package info (click to toggle)
libcudacxx 1.8.1-2
  • links: PTS, VCS
  • area: main
  • in suites: bookworm
  • size: 66,464 kB
  • sloc: cpp: 517,767; ansic: 9,474; python: 6,108; sh: 2,225; asm: 2,154; makefile: 7
file content (37 lines) | stat: -rw-r--r-- 884 bytes parent folder | download
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
---
grand_parent: Synchronization Primitives
parent: cuda::atomic
---

# `cuda::atomic::fetch_max`

Defined in header `<cuda/atomic>`:

```cuda
template <typename T, cuda::thread_scope Scope>
__host__ __device__
T cuda::atomic<T, Scope>::fetch_max(T const& val,
                                    cuda::std::memory_order order
                                      = cuda::std::memory_order_seq_cst);
```

Atomically find the maximum of the value stored in the `cuda::atomic` and `val`.
The maximum is found using [`cuda::std::max`].

## Example

```cuda
#include <cuda/atomic>

__global__ void example_kernel() {
  cuda::atomic<int> a(0);
  auto x = a.fetch_max(1); // Operates as if unsigned.
  auto y = a.load();
  assert(x == 1 && y == 0);
}
```

[See it on Godbolt](https://godbolt.org/z/rexn5T78G){: .btn }


[`cuda::std::max`]: https://en.cppreference.com/w/cpp/algorithm/max