thrust / dependencies /libcudacxx /docs /extended_api /synchronization_primitives /atomic /fetch_max.md
|
Download dependencies/libcudacxx/docs/extended_api/synchronization_primitives/atomic/fetch_max.md from camenduru/thrust: direct link, hf CLI and curl.
- Browser
- Download file 884 Bytes
-
https://huggingface.co/camenduru/thrust/resolve/main/dependencies/libcudacxx/docs/extended_api/synchronization_primitives/atomic/fetch_max.md
- Command line
-
hf download hf://camenduru/thrust/dependencies/libcudacxx/docs/extended_api/synchronization_primitives/atomic/fetch_max.md
-
curl -L -o fetch_max.md https://huggingface.co/camenduru/thrust/resolve/main/dependencies/libcudacxx/docs/extended_api/synchronization_primitives/atomic/fetch_max.md
884 Bytes
metadata
grand_parent: Synchronization Primitives
parent: cuda::atomic
cuda::atomic::fetch_max
Defined in header <cuda/atomic>:
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
#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{: .btn }