Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_GpuParallelReduce.H
Go to the documentation of this file.
1#ifndef AMREX_GPU_PARALLEL_REDUCE_H_
2#define AMREX_GPU_PARALLEL_REDUCE_H_
3#include <AMReX_Config.H>
4
6#include <AMReX_INT.H>
9
10#include <cstddef>
11
12//
13// GPU-aware MPI collectives that operate in place on a Gpu::DeviceVector.
14//
15// These overloads complement the pointer/scalar overloads in
16// AMReX_ParallelReduce.H and AMReX_ParallelDescriptor.H. They live in a
17// separate header (rather than in those low-level headers) because they need
18// the Gpu container/copy machinery (AMReX_GpuContainers.H) that would bloat the
19// ParallelReduce.H headers. This header is also pulled into the AMReX_Gpu.H
20// umbrella for convenience.
21//
22// When AMReX is configured with GPU-aware MPI (ParallelDescriptor::UseGpuAwareMpi())
23// the device buffer is handed to MPI directly, otherwise the data is staged
24// through host (pinned) memory for the collective.
25//
26// Either way the collectives are stream-ordered: work previously launched on
27// Gpu::gpuStream() has completed before MPI reads the vector.
28//
29
30namespace amrex {
31
32namespace ParallelAllReduce {
33
35
36template <typename T>
38{
39 // GPU-unaware case
40#if defined(AMREX_USE_MPI) && defined(AMREX_USE_GPU)
43 Gpu::copy(Gpu::deviceToHost, v.begin(), v.end(), hv.begin());
44 Sum(hv.data(), static_cast<int>(hv.size()), comm);
45 Gpu::copy(Gpu::hostToDevice, hv.begin(), hv.end(), v.begin());
46 return;
47 }
48#endif
49
50 // GPU-aware case
52 Sum(v.data(), static_cast<int>(v.size()), comm);
53}
54
55} // namespace ParallelAllReduce
56
57namespace ParallelReduce {
58
60
61template <typename T>
62void Sum (Gpu::DeviceVector<T>& v, int root, MPI_Comm comm)
63{
64 // GPU-unaware case
65#if defined(AMREX_USE_MPI) && defined(AMREX_USE_GPU)
68 // every rank stages its contribution to host for the reduction
69 Gpu::copy(Gpu::deviceToHost, v.begin(), v.end(), hv.begin());
70 Sum(hv.data(), static_cast<int>(hv.size()), root, comm);
71 // only the root receives the reduced result, so only it copies back
72 if (ParallelDescriptor::MyProc(comm) == root) {
73 Gpu::copy(Gpu::hostToDevice, hv.begin(), hv.end(), v.begin());
74 }
75 return;
76 }
77#endif
78
79 // GPU-aware case
81 Sum(v.data(), static_cast<int>(v.size()), root, comm);
82}
83
84} // namespace ParallelReduce
85
86namespace ParallelDescriptor {
87
89
104template <typename T>
105void Bcast (Gpu::DeviceVector<T>& v, int root, MPI_Comm comm)
106{
107#ifdef AMREX_USE_MPI
108 auto const n = v.size();
109
110#ifdef AMREX_DEBUG
111 // verify the pre-allocation contract (the length broadcast happens on every
112 // rank, so it is collectively safe and cannot deadlock)
113 Long n_root = static_cast<Long>(n);
114 Bcast(&n_root, std::size_t(1), root, comm);
115 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(n_root == static_cast<Long>(n),
116 "ParallelDescriptor::Bcast(Gpu::DeviceVector): receiver must be pre-allocated to the root's length");
117#endif
118
119 // trivial case: 1 rank
120 if (n == 0) { return; }
121
122 // GPU-unaware case
123#ifdef AMREX_USE_GPU
124 if (!UseGpuAwareMpi()) {
126 const bool is_root = (MyProc(comm) == root);
127 // only the root needs to stage its data to host before the broadcast
128 if (is_root) {
129 Gpu::copy(Gpu::deviceToHost, v.begin(), v.end(), hv.begin());
130 }
131 Bcast(hv.data(), static_cast<std::size_t>(n), root, comm);
132 // only the receivers need to copy the broadcast result back to device
133 if (!is_root) {
134 Gpu::copy(Gpu::hostToDevice, hv.begin(), hv.end(), v.begin());
135 }
136 return;
137 }
138#endif
139
140 // GPU-aware case
142 Bcast(v.data(), static_cast<std::size_t>(n), root, comm);
143
144#else // AMREX_USE_MPI
145 amrex::ignore_unused(v, root, comm);
146#endif
147}
148
149} // namespace ParallelDescriptor
150
151} // namespace amrex
152
153#endif /*AMREX_GPU_PARALLEL_REDUCE_H_*/
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:49
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
size_type size() const noexcept
Definition AMReX_PODVector.H:654
iterator begin() noexcept
Definition AMReX_PODVector.H:680
iterator end() noexcept
Definition AMReX_PODVector.H:684
T * data() noexcept
Definition AMReX_PODVector.H:672
amrex_long Long
Definition AMReX_INT.H:30
int MyProc() noexcept
Definition AMReX_ParallelDescriptor.H:128
void Bcast(Gpu::DeviceVector< T > &v, int root, MPI_Comm comm)
Definition AMReX_GpuParallelReduce.H:105
void Sum(Gpu::DeviceVector< T > &v, MPI_Comm comm)
Definition AMReX_GpuParallelReduce.H:37
void Sum(Gpu::DeviceVector< T > &v, int root, MPI_Comm comm)
Definition AMReX_GpuParallelReduce.H:62
void copy(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous st...
Definition AMReX_GpuContainers.H:128
static constexpr DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:106
static constexpr HostToDevice hostToDevice
Definition AMReX_GpuContainers.H:105
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:310
bool UseGpuAwareMpi()
Definition AMReX_ParallelDescriptor.H:113
int MPI_Comm
Definition AMReX_ccse-mpi.H:51
Definition AMReX_Amr.cpp:50
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:259