Skip to content

Commit aab2e8a

Browse files
svalatsvalat
committed
gpu-direct: solve most of the compiler warnings when compiling with CUDA
1 parent 76d46ff commit aab2e8a

5 files changed

Lines changed: 39 additions & 76 deletions

File tree

‎src/fluid/boundary/axis.cpp‎

Lines changed: 7 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ void Axis::ShowConfig() {
2525
void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D<real> Ex1) {
2626
#if DIMENSIONS == 3
2727

28-
idefix::IdefixCommArray1D<real> Ex1Avg = this->Ex1Avg;
28+
auto Ex1Avg = this->Ex1Avg.deviceView();
2929

3030
idefix_for("Ex1_ini",0,data->np_tot[IDIR],
3131
KOKKOS_LAMBDA(int i) {
@@ -40,7 +40,7 @@ void Axis::SymmetrizeEx1Side(int jref, IdefixArray3D<real> Ex1) {
4040
#ifdef WITH_MPI
4141
Kokkos::fence();
4242
// sum along all of the processes on the same r
43-
idefix::MPI_Allreduce(MPI_IN_PLACE, Ex1Avg, data->np_tot[IDIR], realMPI,
43+
idefix::MPI_Allreduce(MPI_IN_PLACE, this->Ex1Avg, data->np_tot[IDIR], realMPI,
4444
MPI_SUM, data->mygrid->AxisComm);
4545
#endif
4646
}
@@ -88,7 +88,8 @@ void Axis::RegularizeCurrentSide(int side) {
8888
jc = data->end[JDIR]-1;
8989
sign = -1;
9090
}
91-
idefix::IdefixCommArray1D<real> BAvg = this->Ex1Avg;
91+
auto BAvg = this->Ex1Avg.deviceView();
92+
auto BAvgComm = this->Ex1Avg;
9293
IdefixArray1D<real> x1 = data->x[IDIR];
9394
IdefixArray1D<real> dx3 = data->dx[KDIR];
9495
IdefixArray1D<real> dx2 = data->dx[JDIR];
@@ -107,7 +108,7 @@ void Axis::RegularizeCurrentSide(int side) {
107108
#ifdef WITH_MPI
108109
Kokkos::fence();
109110
// sum along all of the processes on the same r
110-
MPI_Allreduce(MPI_IN_PLACE, BAvg, data->np_tot[IDIR], realMPI,
111+
MPI_Allreduce(MPI_IN_PLACE, BAvgComm, data->np_tot[IDIR], realMPI,
111112
MPI_SUM, data->mygrid->AxisComm);
112113
#endif
113114
}
@@ -169,7 +170,7 @@ void Axis::FixBx2sAxis(int side) {
169170
// Compute the values of Bx and By that are consistent with BX2 along the axis
170171
#if DIMENSIONS == 3
171172
IdefixArray4D<real> Vs = this->Vs;
172-
idefix::IdefixCommArray2D<real> BAvg = this->BAvg;
173+
auto BAvg = this->BAvg.deviceView();
173174
IdefixArray1D<real> phi = data->x[KDIR];
174175

175176
int jin = 0;
@@ -208,7 +209,7 @@ void Axis::FixBx2sAxis(int side) {
208209
Kokkos::fence();
209210
#ifdef WITH_MPI
210211
// sum along all of the processes on the same r
211-
idefix::MPI_Allreduce(MPI_IN_PLACE, BAvg, 2*data->np_tot[IDIR], realMPI,
212+
idefix::MPI_Allreduce(MPI_IN_PLACE, this->BAvg, 2*data->np_tot[IDIR], realMPI,
212213
MPI_SUM, data->mygrid->AxisComm);
213214
#endif
214215
}

‎src/fluid/constrainedTransport/enforceEMFBoundary.hpp‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -108,9 +108,9 @@ void ConstrainedTransport<Phys>::SymmetrizeEMFShearingBox() {
108108
idfx::pushRegion("Emf::EnforceEMFBoundary");
109109
#if MHD == YES
110110

111-
auto sbEyL = this->sbEyL;
112-
auto sbEyR = this->sbEyR;
113-
auto sbEyRL = this->sbEyRL;
111+
auto sbEyL = this->sbEyL.deviceView();
112+
auto sbEyR = this->sbEyR.deviceView();
113+
auto sbEyRL = this->sbEyRL.deviceView();
114114

115115
IdefixArray3D<real> ey = this->ey;
116116

@@ -142,16 +142,16 @@ void ConstrainedTransport<Phys>::SymmetrizeEMFShearingBox() {
142142
// We send to our left (which, by periodicity, is the right end of the domain)
143143
// our value of sbEyL and get
144144
Kokkos::fence();
145-
idefix::MPI_Sendrecv(sbEyL, size, realMPI, procLeft, 2001,
146-
sbEyR, size, realMPI, procLeft, 2002,
145+
idefix::MPI_Sendrecv(this->sbEyL, size, realMPI, procLeft, 2001,
146+
this->sbEyR, size, realMPI, procLeft, 2002,
147147
data->mygrid->CartComm, &status );
148148
}
149149
if(data->rbound[IDIR]==shearingbox) {
150150
Kokkos::fence();
151151
// We send to our right (which, by periodicity, is the left end (=beginning)
152152
// of the domain) our value of sbEyR and get sbEyL
153-
idefix::MPI_Sendrecv(sbEyR, size, realMPI, procRight, 2002,
154-
sbEyL, size, realMPI, procRight, 2001,
153+
idefix::MPI_Sendrecv(this->sbEyR, size, realMPI, procRight, 2002,
154+
this->sbEyL, size, realMPI, procRight, 2001,
155155
data->mygrid->CartComm, &status );
156156
}
157157
}

‎src/mpi/buffer.hpp‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -86,7 +86,7 @@ class Buffer {
8686
const int kend = box[KDIR][1];
8787
const int offset = this->pointer;
8888

89-
auto arr = this->array;
89+
auto arr = this->array.deviceView();
9090
idefix_for("LoadBuffer4D_var",kbeg,kend,jbeg,jend,ibeg,iend,
9191
KOKKOS_LAMBDA (int k, int j, int i) {
9292
arr(i-ibeg + (j-jbeg)*ni + (k-kbeg)*ninj + offset ) = in(var, k,j,i);
@@ -109,7 +109,7 @@ class Buffer {
109109
const int jend = box[JDIR][1];
110110
const int kend = box[KDIR][1];
111111
const int offset = this->pointer;
112-
auto arr = this->array;
112+
auto arr = this->array.deviceView();
113113

114114
idefix_for("LoadBuffer4D_map",0,map.size(),
115115
kbeg,kend,
@@ -135,7 +135,7 @@ class Buffer {
135135
const int jend = box[JDIR][1];
136136
const int kend = box[KDIR][1];
137137
const int offset = this->pointer;
138-
auto arr = this->array;
138+
auto arr = this->array.deviceView();
139139

140140
idefix_for("UnLoadBuffer3D",kbeg,kend,jbeg,jend,ibeg,iend,
141141
KOKKOS_LAMBDA (int k, int j, int i) {
@@ -160,7 +160,7 @@ class Buffer {
160160
const int kend = box[KDIR][1];
161161
const int offset = this->pointer;
162162

163-
auto arr = this->array;
163+
auto arr = this->array.deviceView();
164164
idefix_for("UnLoadBuffer4D_var",kbeg,kend,jbeg,jend,ibeg,iend,
165165
KOKKOS_LAMBDA (int k, int j, int i) {
166166
out(var,k,j,i) = arr(i-ibeg + (j-jbeg)*ni + (k-kbeg)*ninj + offset );
@@ -185,7 +185,7 @@ class Buffer {
185185
const int kend = box[KDIR][1];
186186
const int offset = this->pointer;
187187

188-
auto arr = this->array;
188+
auto arr = this->array.deviceView();
189189
idefix_for("UnLoadBuffer4D_var_sym",kbeg,kend,jbeg,jend,ibeg,iend,
190190
KOKKOS_LAMBDA (int k, int j, int i) {
191191
const int jinverted = jend-(j-jbeg)-1;
@@ -211,7 +211,7 @@ class Buffer {
211211
const int kend = box[KDIR][1];
212212
const int offset = this->pointer;
213213

214-
auto arr = this->array;
214+
auto arr = this->array.deviceView();
215215
idefix_for("UnLoadBuffer4D_map",0,map.size(),
216216
kbeg,kend,
217217
jbeg,jend,
@@ -239,7 +239,7 @@ class Buffer {
239239
const int kend = box[KDIR][1];
240240
const int offset = this->pointer;
241241

242-
auto arr = this->array;
242+
auto arr = this->array.deviceView();
243243
idefix_for("UnLoadBuffer4D_map_sym",0,map.size(),
244244
kbeg,kend,
245245
jbeg,jend,

‎src/mpi/mpiView.hpp‎

Lines changed: 14 additions & 52 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,6 @@
1616
#include <cassert>
1717
#include <vector>
1818
#include <array>
19-
#include <utility>
2019

2120
#include "idefix.hpp"
2221
#include "buffer.hpp"
@@ -34,25 +33,11 @@ class IdefixCommArrayGpuDirect : public T {
3433
//inherit constructors
3534
using T::T;
3635

37-
IdefixCommArrayGpuDirect(const IdefixCommArrayGpuDirect & orig) {
38-
T & tmp = *this;
39-
tmp = (T&)orig;
40-
}
41-
42-
IdefixCommArrayGpuDirect(IdefixCommArrayGpuDirect && orig) {
43-
T & tmp = *this;
44-
tmp = std::move((T&)orig);
45-
}
46-
47-
IdefixCommArrayGpuDirect<T> & operator=(const IdefixCommArrayGpuDirect & orig) {
48-
T & tmp = *this;
49-
tmp = (T&)orig;
50-
return *this;
51-
}
52-
53-
IdefixCommArrayGpuDirect<T> & operator=(IdefixCommArrayGpuDirect && orig) {
54-
T & tmp = *this;
55-
tmp = std::move((T&)orig);
36+
/**
37+
* Accès it as a simple device view for usiage in kokkos kernels
38+
* (because need to not carry the host part).
39+
*/
40+
T & deviceView(void) {
5641
return *this;
5742
}
5843

@@ -100,33 +85,14 @@ class IdefixCommArrayNoGpuDirect : public T {
10085
//inherit constructors
10186
using T::T;
10287

103-
IdefixCommArrayNoGpuDirect(const IdefixCommArrayNoGpuDirect & orig) {
104-
T & tmp = *this;
105-
tmp = (T&)orig;
106-
this->commArray = orig.commArray;
107-
}
108-
109-
IdefixCommArrayNoGpuDirect(IdefixCommArrayNoGpuDirect && orig) {
110-
T & tmp = *this;
111-
tmp = std::move((T&)orig);
112-
this->commArray = std::move(orig.commArray);
113-
}
114-
115-
IdefixCommArrayNoGpuDirect<T> & operator=(IdefixCommArrayNoGpuDirect & orig) {
116-
T & tmp = *this;
117-
tmp = (T&)orig;
118-
this->commArray = orig.commArray;
119-
return *this;
120-
}
121-
122-
IdefixCommArrayNoGpuDirect<T> & operator=(IdefixCommArrayNoGpuDirect && orig) {
123-
T & tmp = *this;
124-
tmp = std::move((T&)orig);
125-
this->commArray = std::move(orig.commArray);
88+
/**
89+
* Accès it as a simple device view for usiage in kokkos kernels
90+
* (because need to not carry the host part).
91+
*/
92+
T & deviceView(void) {
12693
return *this;
12794
}
12895

129-
13096
/**
13197
* If needed, transfers the data from the device to the host to be ready to make
13298
* a communication.
@@ -223,8 +189,8 @@ template <class T> using MPI_Request_4D = Idefix_MPI_Request< IdefixArray4D<T> >
223189
* to the host if required (when WITH_MPI_GPU_DIRECT is disabled).
224190
*/
225191
template <class T>
226-
int MPI_Sendrecv(IdefixCommArray<T> sendbuf, int sendcount, MPI_Datatype sendtype,
227-
int dest, int sendtag, IdefixCommArray<T> recvbuf, int recvcount,
192+
int MPI_Sendrecv(IdefixCommArray<T> & sendbuf, int sendcount, MPI_Datatype sendtype,
193+
int dest, int sendtag, IdefixCommArray<T> &recvbuf, int recvcount,
228194
MPI_Datatype recvtype, int source, int recvtag, MPI_Comm comm, MPI_Status *status) {
229195
//check size
230196
assert(sendbuf.span() == sendcount);
@@ -233,19 +199,15 @@ int MPI_Sendrecv(IdefixCommArray<T> sendbuf, int sendcount, MPI_Datatype sendtyp
233199
assert(recvbuf.span_is_contiguous());
234200

235201
//make a local allocation if needed
236-
#ifndef WITH_MPI_GPU_DIRECT
237-
sendbuf.syncCommData();
238-
#endif
202+
sendbuf.syncCommData();
239203

240204
//make the communication
241205
const int result = ::MPI_Sendrecv(sendbuf.commData(), sendcount, sendtype, dest, sendtag,
242206
recvbuf.commData(), recvcount, recvtype, source, recvtag,
243207
comm, status);
244208

245209
//copy back if needed
246-
#ifndef WITH_MPI_GPU_DIRECT
247-
recvbuf.syncDeviceData();
248-
#endif
210+
recvbuf.syncDeviceData();
249211

250212
//ok
251213
return result;

‎src/utils/column.cpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -79,7 +79,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
7979
auto column = this->ColumnArray;
8080
auto dV = this->Volume;
8181
auto A = this->Area;
82-
auto localSum = this->localSum;
82+
auto localSum = this->localSum.deviceView();
8383

8484
if(direction==IDIR) {
8585
// Inspired from loop.hpp
@@ -134,7 +134,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
134134
MPI_Status status;
135135
// Get the cumulative sum from previous processes
136136
Kokkos::fence();
137-
idefix::MPI_Recv(localSum, size, realMPI, src, 20, ColumnComm, &status);
137+
idefix::MPI_Recv(this->localSum, size, realMPI, src, 20, ColumnComm, &status);
138138
// Add this to our cumulative sum
139139
idefix_for("Addsum",kb,ke,jb,je,ib,ie,
140140
KOKKOS_LAMBDA(int k, int j, int i) {
@@ -165,7 +165,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
165165
}
166166
// And send it
167167
Kokkos::fence();
168-
idefix::MPI_Send(localSum,size, realMPI, dst, 20, ColumnComm);
168+
idefix::MPI_Send(this->localSum,size, realMPI, dst, 20, ColumnComm);
169169
} // MPIrank small enough
170170
#endif
171171
// If we need it backwards
@@ -195,7 +195,7 @@ void Column::ComputeColumn(IdefixArray4D<real> in, const int var) {
195195

196196
#ifdef WITH_MPI
197197
Kokkos::fence();
198-
idefix::MPI_Bcast(localSum,size, realMPI, MPIsize-1, ColumnComm);
198+
idefix::MPI_Bcast(this->localSum,size, realMPI, MPIsize-1, ColumnComm);
199199
#endif
200200
// All substract the local column from the full column
201201

0 commit comments

Comments
 (0)