Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_GpuLaunchMacrosG.nolint.H
Go to the documentation of this file.
1
// Do not include this header anywhere other than AMReX_GpuLaunchMacrosG.H.
2
// The purpose of this file is to avoid clang-tidy.
3
4
#ifdef AMREX_USE_SYCL
5
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE(TN,TI,block) \
6
{ auto const& amrex_i_tn = TN; \
7
if (!amrex::isEmpty(amrex_i_tn)) { \
8
if (amrex::Gpu::inLaunchRegion()) \
9
{ \
10
const auto amrex_i_ec = amrex::Gpu::ExecutionConfig(amrex_i_tn); \
11
int amrex_i_nthreads_per_block = amrex_i_ec.numThreads.x; \
12
int amrex_i_nthreads_total = amrex_i_nthreads_per_block * amrex_i_ec.numBlocks.x; \
13
auto& amrex_i_q = amrex::Gpu::Device::streamQueue(); \
14
try { \
15
amrex_i_q.submit([&] (sycl::handler& amrex_i_h) { \
16
amrex_i_h.parallel_for(sycl::nd_range<1>(sycl::range<1>(amrex_i_nthreads_total), \
17
sycl::range<1>(amrex_i_nthreads_per_block)), \
18
[=] (sycl::nd_item<1> amrex_i_item) \
19
[[sycl::reqd_work_group_size(AMREX_GPU_MAX_THREADS)]] \
20
[[sycl::reqd_sub_group_size(Gpu::Device::warp_size)]] \
21
{ \
22
for (auto const TI : amrex::Gpu::Range(amrex_i_tn,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
23
block \
24
} \
25
}); \
26
}); \
27
} catch (sycl::exception const& ex) { \
28
amrex::Abort(std::string("LAUNCH: ")+ex.what()+"!!!!!"); \
29
} \
30
} \
31
else { \
32
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
33
block \
34
} \
35
}}}
36
37
#else
38
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE(TN,TI,block) \
39
{ auto const& amrex_i_tn = TN; \
40
if (!amrex::isEmpty(amrex_i_tn)) { \
41
if (amrex::Gpu::inLaunchRegion()) \
42
{ \
43
const auto amrex_i_ec = amrex::Gpu::ExecutionConfig(amrex_i_tn); \
44
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_ec.numBlocks, amrex_i_ec.numThreads, amrex_i_ec.sharedMem, amrex::Gpu::gpuStream(), \
45
[=] AMREX_GPU_DEVICE () noexcept { \
46
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
47
block \
48
} \
49
}); \
50
AMREX_GPU_ERROR_CHECK(); \
51
} \
52
else { \
53
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
54
block \
55
} \
56
}}}
57
#endif
58
59
// two fused launches
60
#ifdef AMREX_USE_SYCL
61
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_2(TN1,TI1,block1,TN2,TI2,block2) \
62
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; \
63
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2)) { \
64
if (amrex::Gpu::inLaunchRegion()) \
65
{ \
66
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
67
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
68
dim3 amrex_i_nblocks = amrex::max(amrex_i_ec1.numBlocks.x, \
69
amrex_i_ec2.numBlocks.x); \
70
amrex_i_nblocks.y = 2; \
71
int amrex_i_nthreads_per_block = amrex_i_ec1.numThreads.x; \
72
int amrex_i_nthreads_total = amrex_i_nthreads_per_block * amrex_i_nblocks.x; \
73
auto& amrex_i_q = amrex::Gpu::Device::streamQueue(); \
74
try { \
75
amrex_i_q.submit([&] (sycl::handler& amrex_i_h) { \
76
amrex_i_h.parallel_for(sycl::nd_range<2>(sycl::range<2>(amrex_i_nthreads_total,2), \
77
sycl::range<2>(amrex_i_nthreads_per_block,1)), \
78
[=] (sycl::nd_item<2> amrex_i_item) \
79
[[sycl::reqd_sub_group_size(Gpu::Device::warp_size)]] \
80
{ \
81
switch (amrex_i_item.get_group(1)) { \
82
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
83
block1 \
84
} \
85
break; \
86
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
87
block2 \
88
} \
89
} \
90
}); \
91
}); \
92
} catch (sycl::exception const& ex) { \
93
amrex::Abort(std::string("LAUNCH: ")+ex.what()+"!!!!!"); \
94
} \
95
} \
96
else { \
97
for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
98
block1 \
99
} \
100
for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
101
block2 \
102
} \
103
}}}
104
105
#else
106
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_2(TN1,TI1,block1,TN2,TI2,block2) \
107
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; \
108
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2)) { \
109
if (amrex::Gpu::inLaunchRegion()) \
110
{ \
111
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
112
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
113
dim3 amrex_i_nblocks = amrex::max(amrex_i_ec1.numBlocks.x, \
114
amrex_i_ec2.numBlocks.x); \
115
amrex_i_nblocks.y = 2; \
116
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_nblocks, amrex_i_ec1.numThreads, 0, amrex::Gpu::gpuStream(), \
117
[=] AMREX_GPU_DEVICE () noexcept { \
118
switch (blockIdx.y) { \
119
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
120
block1 \
121
} \
122
break; \
123
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
124
block2 \
125
} \
126
} \
127
}); \
128
AMREX_GPU_ERROR_CHECK(); \
129
} \
130
else { \
131
for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
132
block1 \
133
} \
134
for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
135
block2 \
136
} \
137
}}}
138
#endif
139
140
// three fused launches
141
#ifdef AMREX_USE_SYCL
142
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_3(TN1,TI1,block1,TN2,TI2,block2,TN3,TI3,block3) \
143
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; auto const& amrex_i_tn3 = TN3; \
144
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2) || !amrex::isEmpty(amrex_i_tn3)) { \
145
if (amrex::Gpu::inLaunchRegion()) \
146
{ \
147
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
148
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
149
const auto amrex_i_ec3 = amrex::Gpu::ExecutionConfig(amrex_i_tn3); \
150
dim3 amrex_i_nblocks = amrex::max(amrex::max(amrex_i_ec1.numBlocks.x, \
151
amrex_i_ec2.numBlocks.x), \
152
amrex_i_ec3.numBlocks.x); \
153
amrex_i_nblocks.y = 3; \
154
int amrex_i_nthreads_per_block = amrex_i_ec1.numThreads.x; \
155
int amrex_i_nthreads_total = amrex_i_nthreads_per_block * amrex_i_nblocks.x; \
156
auto& amrex_i_q = amrex::Gpu::Device::streamQueue(); \
157
try { \
158
amrex_i_q.submit([&] (sycl::handler& amrex_i_h) { \
159
amrex_i_h.parallel_for(sycl::nd_range<2>(sycl::range<2>(amrex_i_nthreads_total,3), \
160
sycl::range<2>(amrex_i_nthreads_per_block,1)), \
161
[=] (sycl::nd_item<2> amrex_i_item) \
162
[[sycl::reqd_sub_group_size(Gpu::Device::warp_size)]] \
163
{ \
164
switch (amrex_i_item.get_group(1)) { \
165
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
166
block1 \
167
} \
168
break; \
169
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
170
block2 \
171
} \
172
break; \
173
case 2: for (auto const TI3 : amrex::Gpu::Range(amrex_i_tn3,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
174
block3 \
175
} \
176
} \
177
}); \
178
}); \
179
} catch (sycl::exception const& ex) { \
180
amrex::Abort(std::string("LAUNCH: ")+ex.what()+"!!!!!"); \
181
} \
182
} \
183
else { \
184
for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
185
block1 \
186
} \
187
for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
188
block2 \
189
} \
190
for (auto const TI3 : amrex::Gpu::Range(amrex_i_tn3)) { \
191
block3 \
192
} \
193
}}}
194
195
#else
196
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_3(TN1,TI1,block1,TN2,TI2,block2,TN3,TI3,block3) \
197
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; auto const& amrex_i_tn3 = TN3; \
198
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2) || !amrex::isEmpty(amrex_i_tn3)) { \
199
if (amrex::Gpu::inLaunchRegion()) \
200
{ \
201
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
202
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
203
const auto amrex_i_ec3 = amrex::Gpu::ExecutionConfig(amrex_i_tn3); \
204
dim3 amrex_i_nblocks = amrex::max(amrex::max(amrex_i_ec1.numBlocks.x, \
205
amrex_i_ec2.numBlocks.x), \
206
amrex_i_ec3.numBlocks.x); \
207
amrex_i_nblocks.y = 3; \
208
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_nblocks, amrex_i_ec1.numThreads, 0, amrex::Gpu::gpuStream(), \
209
[=] AMREX_GPU_DEVICE () noexcept { \
210
switch (blockIdx.y) { \
211
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
212
block1 \
213
} \
214
break; \
215
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
216
block2 \
217
} \
218
break; \
219
case 2: for (auto const TI3 : amrex::Gpu::Range(amrex_i_tn3)) { \
220
block3 \
221
} \
222
} \
223
}); \
224
AMREX_GPU_ERROR_CHECK(); \
225
} \
226
else { \
227
for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
228
block1 \
229
} \
230
for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
231
block2 \
232
} \
233
for (auto const TI3 : amrex::Gpu::Range(amrex_i_tn3)) { \
234
block3 \
235
} \
236
}}}
237
#endif
238
239
#ifdef AMREX_USE_SYCL
240
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE(TN,TI,block) \
241
{ auto const& amrex_i_tn = TN; \
242
if (!amrex::isEmpty(amrex_i_tn)) { \
243
if (amrex::Gpu::inLaunchRegion()) \
244
{ \
245
auto amrex_i_ec = amrex::Gpu::ExecutionConfig(amrex_i_tn); \
246
int amrex_i_nthreads_per_block = amrex_i_ec.numThreads.x; \
247
int amrex_i_nthreads_total = amrex_i_nthreads_per_block * amrex_i_ec.numBlocks.x; \
248
auto& amrex_i_q = amrex::Gpu::Device::streamQueue(); \
249
try { \
250
amrex_i_q.submit([&] (sycl::handler& amrex_i_h) { \
251
amrex_i_h.parallel_for(sycl::nd_range<1>(sycl::range<1>(amrex_i_nthreads_total), \
252
sycl::range<1>(amrex_i_nthreads_per_block)), \
253
[=] (sycl::nd_item<1> amrex_i_item) \
254
[[sycl::reqd_work_group_size(AMREX_GPU_MAX_THREADS)]] \
255
[[sycl::reqd_sub_group_size(Gpu::Device::warp_size)]] \
256
{ \
257
for (auto const TI : amrex::Gpu::Range(amrex_i_tn,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
258
block \
259
} \
260
}); \
261
}); \
262
} catch (sycl::exception const& ex) { \
263
amrex::Abort(std::string("LAUNCH: ")+ex.what()+"!!!!!"); \
264
} \
265
} \
266
else { \
267
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE: cannot call device function from host"); \
268
}}}
269
#else
270
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE(TN,TI,block) \
271
{ auto const& amrex_i_tn = TN; \
272
if (!amrex::isEmpty(amrex_i_tn)) { \
273
if (amrex::Gpu::inLaunchRegion()) \
274
{ \
275
auto amrex_i_ec = amrex::Gpu::ExecutionConfig(amrex_i_tn); \
276
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_ec.numBlocks, amrex_i_ec.numThreads, amrex_i_ec.sharedMem, amrex::Gpu::gpuStream(), \
277
[=] AMREX_GPU_DEVICE () noexcept { \
278
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
279
block \
280
} \
281
}); \
282
AMREX_GPU_ERROR_CHECK(); \
283
} \
284
else { \
285
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE: cannot call device function from host"); \
286
}}}
287
#endif
288
289
// two fused launches
290
#ifdef AMREX_USE_SYCL
291
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2(TN1,TI1,block1,TN2,TI2,block2) \
292
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; \
293
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2)) { \
294
if (amrex::Gpu::inLaunchRegion()) \
295
{ \
296
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
297
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
298
dim3 amrex_i_nblocks = amrex::max(amrex_i_ec1.numBlocks.x, \
299
amrex_i_ec2.numBlocks.x); \
300
amrex_i_nblocks.y = 2; \
301
int amrex_i_nthreads_per_block = amrex_i_ec1.numThreads.x; \
302
int amrex_i_nthreads_total = amrex_i_nthreads_per_block * amrex_i_nblocks.x; \
303
auto& amrex_i_q = amrex::Gpu::Device::streamQueue(); \
304
try { \
305
amrex_i_q.submit([&] (sycl::handler& amrex_i_h) { \
306
amrex_i_h.parallel_for(sycl::nd_range<2>(sycl::range<2>(amrex_i_nthreads_total,2), \
307
sycl::range<2>(amrex_i_nthreads_per_block,1)), \
308
[=] (sycl::nd_item<2> amrex_i_item) \
309
[[sycl::reqd_sub_group_size(Gpu::Device::warp_size)]] \
310
{ \
311
switch (amrex_i_item.get_group(1)) { \
312
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
313
block1 \
314
} \
315
break; \
316
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
317
block2 \
318
} \
319
} \
320
}); \
321
}); \
322
} catch (sycl::exception const& ex) { \
323
amrex::Abort(std::string("LAUNCH: ")+ex.what()+"!!!!!"); \
324
} \
325
} \
326
else { \
327
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2: cannot call device function from host"); \
328
}}}
329
#else
330
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2(TN1,TI1,block1,TN2,TI2,block2) \
331
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; \
332
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2)) { \
333
if (amrex::Gpu::inLaunchRegion()) \
334
{ \
335
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
336
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
337
dim3 amrex_i_nblocks = amrex::max(amrex_i_ec1.numBlocks.x, \
338
amrex_i_ec2.numBlocks.x); \
339
amrex_i_nblocks.y = 2; \
340
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_nblocks, amrex_i_ec1.numThreads, 0, amrex::Gpu::gpuStream(), \
341
[=] AMREX_GPU_DEVICE () noexcept { \
342
switch (blockIdx.y) { \
343
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
344
block1 \
345
} \
346
break; \
347
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
348
block2 \
349
} \
350
} \
351
}); \
352
AMREX_GPU_ERROR_CHECK(); \
353
} \
354
else { \
355
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2: cannot call device function from host"); \
356
}}}
357
#endif
358
359
// three fused launches
360
#ifdef AMREX_USE_SYCL
361
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3(TN1,TI1,block1,TN2,TI2,block2,TN3,TI3,block3) \
362
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; auto const& amrex_i_tn3 = TN3; \
363
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2) || !amrex::isEmpty(amrex_i_tn3)) { \
364
if (amrex::Gpu::inLaunchRegion()) \
365
{ \
366
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
367
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
368
const auto amrex_i_ec3 = amrex::Gpu::ExecutionConfig(amrex_i_tn3); \
369
dim3 amrex_i_nblocks = amrex::max(amrex::max(amrex_i_ec1.numBlocks.x, \
370
amrex_i_ec2.numBlocks.x), \
371
amrex_i_ec3.numBlocks.x); \
372
amrex_i_nblocks.y = 3; \
373
int amrex_i_nthreads_per_block = amrex_i_ec1.numThreads.x; \
374
int amrex_i_nthreads_total = amrex_i_nthreads_per_block * amrex_i_nblocks.x; \
375
auto& amrex_i_q = amrex::Gpu::Device::streamQueue(); \
376
try { \
377
amrex_i_q.submit([&] (sycl::handler& amrex_i_h) { \
378
amrex_i_h.parallel_for(sycl::nd_range<2>(sycl::range<2>(amrex_i_nthreads_total,3), \
379
sycl::range<2>(amrex_i_nthreads_per_block,1)), \
380
[=] (sycl::nd_item<2> amrex_i_item) \
381
[[sycl::reqd_sub_group_size(Gpu::Device::warp_size)]] \
382
{ \
383
switch (amrex_i_item.get_group(1)) { \
384
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
385
block1 \
386
} \
387
break; \
388
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
389
block2 \
390
} \
391
break; \
392
case 2: for (auto const TI3 : amrex::Gpu::Range(amrex_i_tn3,amrex_i_item.get_global_id(0),amrex_i_item.get_global_range(0))) { \
393
block3 \
394
} \
395
} \
396
}); \
397
}); \
398
} catch (sycl::exception const& ex) { \
399
amrex::Abort(std::string("LAUNCH: ")+ex.what()+"!!!!!"); \
400
} \
401
} \
402
else { \
403
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3: cannot call device function from host"); \
404
}}}
405
#else
406
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3(TN1,TI1,block1,TN2,TI2,block2,TN3,TI3,block3) \
407
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; auto const& amrex_i_tn3 = TN3; \
408
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2) || !amrex::isEmpty(amrex_i_tn3)) { \
409
if (amrex::Gpu::inLaunchRegion()) \
410
{ \
411
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
412
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
413
const auto amrex_i_ec3 = amrex::Gpu::ExecutionConfig(amrex_i_tn3); \
414
dim3 amrex_i_nblocks = amrex::max(amrex::max(amrex_i_ec1.numBlocks.x, \
415
amrex_i_ec2.numBlocks.x), \
416
amrex_i_ec3.numBlocks.x); \
417
amrex_i_nblocks.y = 3; \
418
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_nblocks, amrex_i_ec1.numThreads, 0, amrex::Gpu::gpuStream(), \
419
[=] AMREX_GPU_DEVICE () noexcept { \
420
switch (blockIdx.y) { \
421
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
422
block1 \
423
} \
424
break; \
425
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
426
block2 \
427
} \
428
break; \
429
case 2: for (auto const TI3 : amrex::Gpu::Range(amrex_i_tn3)) { \
430
block3 \
431
} \
432
} \
433
}); \
434
AMREX_GPU_ERROR_CHECK(); \
435
} \
436
else { \
437
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3: cannot call device function from host"); \
438
}}}
439
#endif
440
441
// FOR_1D
442
443
#define AMREX_HOST_DEVICE_FOR_1D(n,i,block) \
444
{ \
445
auto const& amrex_i_n = n; \
446
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
447
if (amrex::Gpu::inLaunchRegion()) { \
448
amrex::ParallelFor(amrex_i_n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
449
} else { \
450
auto amrex_i_lambda = [=] (amrex_i_inttype i) noexcept block; \
451
AMREX_PRAGMA_SIMD \
452
for (amrex_i_inttype i = 0; i < amrex_i_n; ++i) amrex_i_lambda(i); \
453
} \
454
}
455
456
#define AMREX_HOST_DEVICE_PARALLEL_FOR_1D(n,i,block) \
457
{ \
458
auto const& amrex_i_n = n; \
459
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
460
if (amrex::Gpu::inLaunchRegion()) { \
461
amrex::ParallelFor(amrex_i_n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
462
} else { \
463
auto amrex_i_lambda = [=] (amrex_i_inttype i) noexcept block; \
464
AMREX_PRAGMA_SIMD \
465
for (amrex_i_inttype i = 0; i < amrex_i_n; ++i) amrex_i_lambda(i); \
466
} \
467
}
468
469
#define AMREX_FOR_1D(n,i,block) \
470
{ \
471
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
472
amrex::ParallelFor(n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
473
}
474
475
#define AMREX_PARALLEL_FOR_1D(n,i,block) \
476
{ \
477
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
478
amrex::ParallelFor(n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
479
}
480
481
// FOR_3D
482
483
#define AMREX_HOST_DEVICE_FOR_3D(box,i,j,k,block) \
484
{ \
485
auto const& amrex_i_box = box; \
486
if (amrex::Gpu::inLaunchRegion()) { \
487
amrex::ParallelFor(amrex_i_box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
488
} else { \
489
amrex::LoopConcurrentOnCpu(amrex_i_box,[=] (int i, int j, int k) noexcept block); \
490
} \
491
}
492
493
#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D(box,i,j,k,block) \
494
{ \
495
auto const& amrex_i_box = box; \
496
if (amrex::Gpu::inLaunchRegion()) { \
497
amrex::ParallelFor(amrex_i_box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
498
} else { \
499
amrex::LoopConcurrentOnCpu(amrex_i_box,[=] (int i, int j, int k) noexcept block); \
500
} \
501
}
502
503
#define AMREX_FOR_3D(box,i,j,k,block) \
504
{ \
505
amrex::ParallelFor(box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
506
}
507
508
#define AMREX_PARALLEL_FOR_3D(box,i,j,k,block) \
509
{ \
510
amrex::ParallelFor(box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
511
}
512
513
// FOR_4D
514
515
#define AMREX_HOST_DEVICE_FOR_4D(box,ncomp,i,j,k,n,block) \
516
{ \
517
auto const& amrex_i_box = box; \
518
auto const& amrex_i_ncomp = ncomp; \
519
if (amrex::Gpu::inLaunchRegion()) { \
520
amrex::ParallelFor(amrex_i_box,amrex_i_ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
521
} else { \
522
amrex::LoopConcurrentOnCpu(amrex_i_box,amrex_i_ncomp,[=] (int i, int j, int k, int n) noexcept block); \
523
} \
524
}
525
526
#define AMREX_HOST_DEVICE_PARALLEL_FOR_4D(box,ncomp,i,j,k,n,block) \
527
{ \
528
auto const& amrex_i_box = box; \
529
auto const& amrex_i_ncomp = ncomp; \
530
if (amrex::Gpu::inLaunchRegion()) { \
531
amrex::ParallelFor(amrex_i_box,amrex_i_ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
532
} else { \
533
amrex::LoopConcurrentOnCpu(amrex_i_box,amrex_i_ncomp,[=] (int i, int j, int k, int n) noexcept block); \
534
} \
535
}
536
537
#define AMREX_FOR_4D(box,ncomp,i,j,k,n,block) \
538
{ \
539
amrex::ParallelFor(box,ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
540
}
541
542
#define AMREX_PARALLEL_FOR_4D(box,ncomp,i,j,k,n,block) \
543
{ \
544
amrex::ParallelFor(box,ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
545
}
Src
Base
AMReX_GpuLaunchMacrosG.nolint.H
Generated by
1.15.0
1.15.0