Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
math.hip
Go to the documentation of this file.
1/*
2 Copyright (c) 2021-2026, The Neko Authors
3 All rights reserved.
4
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions
7 are met:
8
9 * Redistributions of source code must retain the above copyright
10 notice, this list of conditions and the following disclaimer.
11
12 * Redistributions in binary form must reproduce the above
13 copyright notice, this list of conditions and the following
14 disclaimer in the documentation and/or other materials provided
15 with the distribution.
16
17 * Neither the name of the authors nor the names of its
18 contributors may be used to endorse or promote products derived
19 from this software without specific prior written permission.
20
21 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 POSSIBILITY OF SUCH DAMAGE.
33*/
34
35#include <hip/hip_runtime.h>
36#include <device/device_config.h>
37#include <device/hip/check.h>
38#include <device/hip/buffer.h>
39#include <device/hip/unified.h>
40#include "math_kernel.h"
41
42extern "C" {
43
44#include <math/bcknd/device/device_mpi_reduce.h>
45#include <math/bcknd/device/device_mpi_op.h>
46
47#ifdef HAVE_RCCL
48#include <math/bcknd/device/device_nccl_reduce.h>
49#include <math/bcknd/device/device_nccl_op.h>
50#endif
51
55 void hip_copy(void *a, void *b, int *n, hipStream_t strm) {
56 if (hip_zerocopy()) {
57 if (*n == 0) return;
58 /* The pointers may alias pageable host memory, which hipMemcpy
59 treats as a staged pageable copy (slow, serializing); copy
60 with a kernel instead, which runs at full memory bandwidth
61 on any pointer */
62 const dim3 nthrds(1024, 1, 1);
63 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
64
65 hipLaunchKernelGGL(HIP_KERNEL_NAME(copy_kernel<real>),
66 nblcks, nthrds, 0, strm,
67 (real *) a, (const real *) b, *n);
68 HIP_CHECK(hipGetLastError());
69 return;
70 }
71 HIP_CHECK(hipMemcpyAsync(a, b, (*n) * sizeof(real),
72 hipMemcpyDeviceToDevice, strm));
73 }
74
78 void hip_masked_copy_0(void *a, void *b, void *mask,
79 int *n, int *m, hipStream_t strm) {
80
81 const dim3 nthrds(1024, 1, 1);
82 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
83
84 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_copy_kernel_0<real>),
85 nblcks, nthrds, 0, strm, (real *) a,
86 (real *) b, (int *) mask, *n, *m);
87
88 HIP_CHECK(hipGetLastError());
89
90 }
91
95 void hip_masked_copy_aligned(void *a, void *b, void *mask,
96 int *n, int *m, hipStream_t strm) {
97
98 const dim3 nthrds(1024, 1, 1);
99 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
100
101 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_copy_kernel_aligned<real>),
102 nblcks, nthrds, 0, strm, (real *) a,
103 (real *) b, (int *) mask, *n, *m);
104
105 HIP_CHECK(hipGetLastError());
106
107 }
108
112 void hip_masked_gather_copy(void *a, void *b, void *mask,
113 int *n, int *m, hipStream_t strm) {
114
115 const dim3 nthrds(1024, 1, 1);
116 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
117
118 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_gather_copy_kernel<real>),
119 nblcks, nthrds, 0, strm, (real *) a,
120 (real *) b, (int *) mask, *n, *m);
121 HIP_CHECK(hipGetLastError());
122
123 }
124
128 void hip_masked_gather_copy_aligned(void *a, void *b, void *mask,
129 int *n, int *m, hipStream_t strm) {
130
131 const dim3 nthrds(1024, 1, 1);
132 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
133
134 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_gather_copy_aligned_kernel<real>),
135 nblcks, nthrds, 0, strm, (real *) a,
136 (real *) b, (int *) mask, *n, *m);
137 HIP_CHECK(hipGetLastError());
138
139 }
140
144 void hip_face_masked_gather_copy(void *a, void *b, void *mask,
145 void *facet, int *n1, int *n2, int *lx,
146 int *ly, int *lz, int *m,
147 hipStream_t strm) {
148
149 const dim3 nthrds(1024, 1, 1);
150 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
151
152 hipLaunchKernelGGL(HIP_KERNEL_NAME(face_masked_gather_copy_kernel<real>),
153 nblcks, nthrds, 0, strm, (real *) a, (real *) b,
154 (int *) mask, (int *) facet, *n1, *n2, *lx, *ly, *lz,
155 *m);
156 HIP_CHECK(hipGetLastError());
157
158 }
159
160
164 void hip_masked_scatter_copy(void *a, void *b, void *mask,
165 int *n, int *m, hipStream_t strm) {
166
167 const dim3 nthrds(1024, 1, 1);
168 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
169
170 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_scatter_copy_kernel<real>),
171 nblcks, nthrds, 0, strm, (real *) a,
172 (real *) b, (int *) mask, *n, *m);
173
174 HIP_CHECK(hipGetLastError());
175
176 }
177
181 void hip_masked_scatter_copy_aligned(void *a, void *b, void *mask,
182 int *n, int *m, hipStream_t strm) {
183
184 const dim3 nthrds(1024, 1, 1);
185 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
186
187 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_scatter_copy_aligned_kernel<real>),
188 nblcks, nthrds, 0, strm, (real *) a,
189 (real *) b, (int *) mask, *n, *m);
190
191 HIP_CHECK(hipGetLastError());
192
193 }
194
198 void hip_masked_atomic_reduction(void *a, void *b, void *mask,
199 int *n, int *m, hipStream_t strm) {
200
201 const dim3 nthrds(MASKED_ATOMIC_RED_BLOCK, 1, 1);
202 const dim3 nblcks(((*m) + MASKED_ATOMIC_RED_BLOCK - 1) /
203 MASKED_ATOMIC_RED_BLOCK, 1, 1);
204
205 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_atomic_reduction_kernel<real>),
206 nblcks, nthrds, 0, strm, (real *) a,
207 (real *) b, (int *) mask, *n, *m);
208
209 HIP_CHECK(hipGetLastError());
210
211 }
212
216 void hip_cfill_mask(void* a, real* c, int* size, void* mask, int* mask_size,
217 hipStream_t strm) {
218
219 const dim3 nthrds(1024, 1, 1);
220 const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
221
222 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_mask_kernel<real>),
223 nblcks, nthrds, 0, strm, (real*)a,
224 *c, *size, (int*)mask, *mask_size);
225
226 HIP_CHECK(hipGetLastError());
227 }
228
232 void hip_rzero(void *a, int *n, hipStream_t strm) {
233 if (hip_zerocopy()) {
234 if (*n == 0) return;
235 /* a may alias pageable host memory, which hipMemset rejects;
236 zero with a kernel instead (works on any pointer) */
237 const dim3 nthrds(1024, 1, 1);
238 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
239
240 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_kernel<real>),
241 nblcks, nthrds, 0, strm,
242 (real *) a, (real) 0.0, *n);
243 HIP_CHECK(hipGetLastError());
244 return;
245 }
246 HIP_CHECK(hipMemsetAsync(a, 0, (*n) * sizeof(real), strm));
247 }
248
252 void hip_cmult(void *a, real *c, int *n, hipStream_t strm) {
253
254 const dim3 nthrds(1024, 1, 1);
255 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
256
257 hipLaunchKernelGGL(HIP_KERNEL_NAME(cmult_kernel<real>),
258 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
259 HIP_CHECK(hipGetLastError());
260
261 }
262
266 void hip_cmult2(void *a, void *b, real *c, int *n, hipStream_t strm) {
267
268 const dim3 nthrds(1024, 1, 1);
269 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
270
271 hipLaunchKernelGGL(HIP_KERNEL_NAME(cmult2_kernel<real>),
272 nblcks, nthrds, 0, strm, (real *) a,(real *) b, *c, *n);
273 HIP_CHECK(hipGetLastError());
274
275 }
276
280 void hip_cdiv(void *a, real *c, int *n, hipStream_t strm) {
281
282 const dim3 nthrds(1024, 1, 1);
283 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
284
285 hipLaunchKernelGGL(HIP_KERNEL_NAME(cdiv_kernel<real>),
286 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
287 HIP_CHECK(hipGetLastError());
288
289 }
290
294 void hip_cdiv2(void *a, void *b, real *c, int *n, hipStream_t strm) {
295
296 const dim3 nthrds(1024, 1, 1);
297 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
298
299 hipLaunchKernelGGL(HIP_KERNEL_NAME(cdiv2_kernel<real>),
300 nblcks, nthrds, 0, strm, (real *) a,(real *) b, *c, *n);
301 HIP_CHECK(hipGetLastError());
302
303 }
304
308 void hip_radd(void *a, real *c, int *n, hipStream_t strm) {
309
310 const dim3 nthrds(1024, 1, 1);
311 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
312
313 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd_kernel<real>),
314 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
315 HIP_CHECK(hipGetLastError());
316 }
317
322 void hip_cadd2(void *a, void *b, real *c, int *n, hipStream_t strm) {
323
324 const dim3 nthrds(1024, 1, 1);
325 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
326
327 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd2_kernel<real>),
328 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *c, *n);
329 HIP_CHECK(hipGetLastError());
330 }
331
336 void hip_cwrap(void *a, real *min_val, real *max_val, int *n,
337 hipStream_t strm) {
338
339 const dim3 nthrds(1024, 1, 1);
340 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
341
342 hipLaunchKernelGGL(HIP_KERNEL_NAME(cwrap_kernel<real>),
343 nblcks, nthrds, 0, strm, (real *) a,
344 *min_val, *max_val, *n);
345 HIP_CHECK(hipGetLastError());
346 }
347
351 void hip_sqrt_inplace(void *a, int *n, hipStream_t strm) {
352
353 const dim3 nthrds(1024, 1, 1);
354 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
355
356 hipLaunchKernelGGL(HIP_KERNEL_NAME(sqrt_inplace_kernel<real>),
357 nblcks, nthrds, 0, strm, (real *) a, *n);
358 HIP_CHECK(hipGetLastError());
359 }
360
364 void hip_power(void *ap, void *a, real *p, int *n, hipStream_t strm) {
365
366 const dim3 nthrds(1024, 1, 1);
367 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
368
369 hipLaunchKernelGGL(HIP_KERNEL_NAME(power_kernel<real>),
370 nblcks, nthrds, 0, strm,
371 (real *) ap, (real *) a, *p, *n);
372 HIP_CHECK(hipGetLastError());
373 }
374
378 void hip_cfill(void *a, real *c, int *n, hipStream_t strm) {
379
380 const dim3 nthrds(1024, 1, 1);
381 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
382
383 if (*n > 0) {
384 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_kernel<real>),
385 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
386 HIP_CHECK(hipGetLastError());
387 }
388
389 }
390
395 void hip_add2(void *a, void *b, int *n, hipStream_t strm) {
396
397 const dim3 nthrds(1024, 1, 1);
398 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
399
400 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2_kernel<real>),
401 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *n);
402 HIP_CHECK(hipGetLastError());
403
404 }
405
410 void hip_add3(void *a, void *b, void *c, int *n, hipStream_t strm) {
411
412 const dim3 nthrds(1024, 1, 1);
413 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
414
415 hipLaunchKernelGGL(HIP_KERNEL_NAME(add3_kernel<real>),
416 nblcks, nthrds, 0, strm, (real *) a,
417 (real *) b, (real *) c, *n);
418 HIP_CHECK(hipGetLastError());
419 }
420
425 void hip_add4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm) {
426
427 const dim3 nthrds(1024, 1, 1);
428 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
429
430 hipLaunchKernelGGL(HIP_KERNEL_NAME(add4_kernel<real>),
431 nblcks, nthrds, 0, strm, (real *) a,
432 (real *) b, (real *) c, (real *) d, *n);
433 HIP_CHECK(hipGetLastError());
434 }
435
441 void hip_add2s1(void *a, void *b, real *c1, int *n, hipStream_t strm) {
442
443 const dim3 nthrds(1024, 1, 1);
444 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
445
446 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s1_kernel<real>),
447 nblcks, nthrds, 0, strm, (real *) a,
448 (real *) b, *c1, *n);
449 HIP_CHECK(hipGetLastError());
450 }
451
457 void hip_add2s2(void *a, void *b, real *c1, int *n, hipStream_t strm) {
458
459 const dim3 nthrds(1024, 1, 1);
460 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
461
462 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s2_kernel<real>),
463 nblcks, nthrds, 0, strm, (real *) a, (real *) b,
464 *c1, *n);
465 HIP_CHECK(hipGetLastError());
466 }
467
474 void hip_add2s2_many(void *x, void **p, void *alpha, int *j, int *n,
475 hipStream_t strm) {
476
477 const dim3 nthrds(1024, 1, 1);
478 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
479
480 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s2_many_kernel<real>),
481 nblcks, nthrds, 0, strm, (real *) x,
482 (const real **) p, (real *) alpha, *j, *n);
483 HIP_CHECK(hipGetLastError());
484
485 }
486
492 void hip_addsqr2s2(void *a, void *b, real *c1, int *n, hipStream_t strm) {
493
494 const dim3 nthrds(1024, 1, 1);
495 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
496
497 hipLaunchKernelGGL(HIP_KERNEL_NAME(addsqr2s2_kernel<real>),
498 nblcks, nthrds, 0, strm, (real *) a,
499 (real *) b, *c1, *n);
500 HIP_CHECK(hipGetLastError());
501 }
502
508 void hip_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n,
509 hipStream_t strm) {
510
511 const dim3 nthrds(1024, 1, 1);
512 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
513
514 hipLaunchKernelGGL(HIP_KERNEL_NAME(add3s2_kernel<real>),
515 nblcks, nthrds, 0, strm, (real *) a,
516 (real *) b, (real *) c, *c1, *c2, *n);
517 HIP_CHECK(hipGetLastError());
518 }
519
525 void hip_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2,
526 real *c3, int *n, hipStream_t strm) {
527
528 const dim3 nthrds(1024, 1, 1);
529 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
530
531 hipLaunchKernelGGL(HIP_KERNEL_NAME(add4s3_kernel<real>),
532 nblcks, nthrds, 0, strm, (real *) a,
533 (real *) b, (real *) c, (real *) d, *c1, *c2, *c3, *n);
534 HIP_CHECK(hipGetLastError());
535 }
536
542 void hip_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1,
543 real *c2, real *c3, real *c4, int *n, hipStream_t strm) {
544
545 const dim3 nthrds(1024, 1, 1);
546 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
547
548 hipLaunchKernelGGL(HIP_KERNEL_NAME(add5s4_kernel<real>),
549 nblcks, nthrds, 0, strm, (real *) a,
550 (real *) b, (real *) c, (real *) d, (real *) e,
551 *c1, *c2, *c3, *c4, *n);
552 HIP_CHECK(hipGetLastError());
553 }
554
559 void hip_invcol1(void *a, int *n, hipStream_t strm) {
560
561 const dim3 nthrds(1024, 1, 1);
562 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
563
564 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol1_kernel<real>),
565 nblcks, nthrds, 0, strm, (real *) a, *n);
566 HIP_CHECK(hipGetLastError());
567 }
568
573 void hip_invcol2(void *a, void *b, int *n, hipStream_t strm) {
574
575 const dim3 nthrds(1024, 1, 1);
576 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
577
578 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol2_kernel<real>),
579 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *n);
580 HIP_CHECK(hipGetLastError());
581 }
582
587 void hip_invcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
588
589 const dim3 nthrds(1024, 1, 1);
590 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
591
592 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol3_kernel<real>),
593 nblcks, nthrds, 0, strm, (real *) a,
594 (real *) b, (real *) c, *n);
595 HIP_CHECK(hipGetLastError());
596 }
597
602 void hip_col2(void *a, void *b, int *n, hipStream_t strm) {
603
604 const dim3 nthrds(1024, 1, 1);
605 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
606
607 hipLaunchKernelGGL(HIP_KERNEL_NAME(col2_kernel<real>),
608 nblcks, nthrds, 0, strm, (real *) a,
609 (real *) b, *n);
610 HIP_CHECK(hipGetLastError());
611 }
612
617 void hip_col3(void *a, void *b, void *c, int *n, hipStream_t strm) {
618
619 const dim3 nthrds(1024, 1, 1);
620 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
621
622 hipLaunchKernelGGL(HIP_KERNEL_NAME(col3_kernel<real>),
623 nblcks, nthrds, 0, strm, (real *) a,
624 (real *) b, (real *) c, *n);
625 HIP_CHECK(hipGetLastError());
626 }
627
632 void hip_subcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
633
634 const dim3 nthrds(1024, 1, 1);
635 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
636
637 hipLaunchKernelGGL(HIP_KERNEL_NAME(subcol3_kernel<real>),
638 nblcks, nthrds, 0, strm, (real *) a,
639 (real *) b, (real *) c, *n);
640 HIP_CHECK(hipGetLastError());
641 }
642
647 void hip_sub2(void *a, void *b, int *n, hipStream_t strm) {
648
649 const dim3 nthrds(1024, 1, 1);
650 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
651
652 hipLaunchKernelGGL(HIP_KERNEL_NAME(sub2_kernel<real>),
653 nblcks, nthrds, 0, strm, (real *) a,
654 (real *) b, *n);
655 HIP_CHECK(hipGetLastError());
656 }
657
662 void hip_sub3(void *a, void *b, void *c, int *n, hipStream_t strm) {
663
664 const dim3 nthrds(1024, 1, 1);
665 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
666
667 hipLaunchKernelGGL(HIP_KERNEL_NAME(sub3_kernel<real>),
668 nblcks, nthrds, 0, strm, (real *) a,
669 (real *) b, (real *) c, *n);
670 HIP_CHECK(hipGetLastError());
671 }
672
677 void hip_addcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
678
679 const dim3 nthrds(1024, 1, 1);
680 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
681
682 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol3_kernel<real>),
683 nblcks, nthrds, 0, strm, (real *) a,
684 (real *) b, (real *) c, *n);
685 HIP_CHECK(hipGetLastError());
686 }
687
692 void hip_addcol4(void *a, void *b, void *c, void *d, int *n,
693 hipStream_t strm) {
694
695 const dim3 nthrds(1024, 1, 1);
696 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
697
698 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol4_kernel<real>),
699 nblcks, nthrds, 0, strm, (real *) a,
700 (real *) b, (real *) c, (real *) d, *n);
701 HIP_CHECK(hipGetLastError());
702 }
703
708 void hip_addcol3s2(void *a, void *b, void *c, real *s, int *n,
709 hipStream_t strm) {
710
711 const dim3 nthrds(1024, 1, 1);
712 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
713
714 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol3s2_kernel<real>),
715 nblcks, nthrds, 0, strm, (real *) a,
716 (real *) b, (real *) c, *s, *n);
717 HIP_CHECK(hipGetLastError());
718 }
719
724 void hip_vdot3(void *dot, void *u1, void *u2, void *u3,
725 void *v1, void *v2, void *v3, int *n,
726 hipStream_t strm) {
727
728 const dim3 nthrds(1024, 1, 1);
729 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
730
731 hipLaunchKernelGGL(HIP_KERNEL_NAME(vdot3_kernel<real>),
732 nblcks, nthrds, 0, strm,
733 (real *) dot, (real *) u1, (real *) u2, (real *) u3,
734 (real *) v1, (real *) v2, (real *) v3, *n);
735 HIP_CHECK(hipGetLastError());
736 }
737
742 void hip_vcross(void *u1, void *u2, void *u3,
743 void *v1, void *v2, void *v3,
744 void *w1, void *w2, void *w3,
745 int *n, hipStream_t strm) {
746
747 const dim3 nthrds(1024, 1, 1);
748 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
749
750 hipLaunchKernelGGL(HIP_KERNEL_NAME(vcross_kernel<real>),
751 nblcks, nthrds, 0, strm,
752 (real *) u1, (real *) u2, (real *) u3,
753 (real *) v1, (real *) v2, (real *) v3,
754 (real *) w1, (real *) w2, (real *) w3, *n);
755 HIP_CHECK(hipGetLastError());
756 }
757
758
759 /*
760 * Reduction buffers, owned by the device layer and released
761 * on device teardown (hip_buffer_free_all in hip_finalize)
762 */
763 hip_buffer_t redbuf = HIP_BUFFER_INIT;
764 hip_buffer_t redbuf_xp = HIP_BUFFER_INIT;
765
766 void hip_redbuf_check_alloc(int nb) {
767 hip_buffer_reserve(&redbuf, (nb + 1) * sizeof(real));
768 }
769
770 void hip_redbuf_check_alloc_xp(int nb) {
771 hip_buffer_reserve(&redbuf_xp, (nb + 1) * sizeof(real_xp));
772 }
773
777 void hip_global_reduce_add(real * bufred, void * bufred_d, int n,
778 const hipStream_t stream) {
779 #ifdef HAVE_RCCL
780 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
781 DEVICE_NCCL_SUM, stream);
782 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
783 hipMemcpyDeviceToHost, stream));
784 HIP_CHECK(hipStreamSynchronize(stream));
785 #elif HAVE_DEVICE_MPI
786 HIP_CHECK(hipStreamSynchronize(stream));
787 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_SUM);
788 #else
789 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
790 hipMemcpyDeviceToHost, stream));
791 HIP_CHECK(hipStreamSynchronize(stream));
792 #endif
793 }
794
798 void hip_global_reduce_add_xp(real_xp * bufred, void * bufred_d, int n,
799 const hipStream_t stream) {
800 #ifdef HAVE_RCCL
801 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real_xp),
802 DEVICE_NCCL_SUM, stream);
803 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real_xp),
804 hipMemcpyDeviceToHost, stream));
805 HIP_CHECK(hipStreamSynchronize(stream));
806 #elif HAVE_DEVICE_MPI
807 HIP_CHECK(hipStreamSynchronize(stream));
808 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real_xp), DEVICE_MPI_SUM);
809 #else
810 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real_xp),
811 hipMemcpyDeviceToHost, stream));
812 HIP_CHECK(hipStreamSynchronize(stream));
813 #endif
814 }
815
816
820 void hip_global_reduce_max(real * bufred, void * bufred_d, int n,
821 const hipStream_t stream) {
822 #ifdef HAVE_RCCL
823 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
824 DEVICE_NCCL_MAX, stream);
825 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
826 hipMemcpyDeviceToHost, stream));
827 HIP_CHECK(hipStreamSynchronize(stream));
828 #elif HAVE_DEVICE_MPI
829 HIP_CHECK(hipStreamSynchronize(stream));
830 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MAX);
831 #else
832 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
833 hipMemcpyDeviceToHost, stream));
834 HIP_CHECK(hipStreamSynchronize(stream));
835 #endif
836 }
837
841 void hip_global_reduce_min(real * bufred, void * bufred_d, int n,
842 const hipStream_t stream) {
843 #ifdef HAVE_RCCL
844 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
845 DEVICE_NCCL_MIN, stream);
846 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
847 hipMemcpyDeviceToHost, stream));
848 HIP_CHECK(hipStreamSynchronize(stream));
849 #elif HAVE_DEVICE_MPI
850 HIP_CHECK(hipStreamSynchronize(stream));
851 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MIN);
852 #else
853 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
854 hipMemcpyDeviceToHost, stream));
855 HIP_CHECK(hipStreamSynchronize(stream));
856 #endif
857 }
858
863 real hip_vlsc3(void *u, void *v, void *w, int *n, hipStream_t stream) {
864
865 const dim3 nthrds(1024, 1, 1);
866 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
867 const int nb = ((*n) + 1024 - 1)/ 1024;
868
869 hip_redbuf_check_alloc(nb);
870
871
872 hipLaunchKernelGGL(HIP_KERNEL_NAME(vlsc3_kernel<real>),
873 nblcks, nthrds, 0, stream,
874 (real *) u, (real *) v,
875 (real *) w, (real *) redbuf.dev, *n);
876 HIP_CHECK(hipGetLastError());
877 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
878 1, 1024, 0, stream, (real *) redbuf.dev, nb);
879 HIP_CHECK(hipGetLastError());
880
881 HIP_CHECK(hipMemcpyAsync(((real *) redbuf.host), (real *) redbuf.dev,
882 sizeof(real), hipMemcpyDeviceToHost, stream));
883 HIP_CHECK(hipStreamSynchronize(stream));
884
885 return ((real *) redbuf.host)[0];
886 }
887
892 real_xp hip_glsc3(void *a, void *b, void *c, int *n, hipStream_t stream) {
893
894 const dim3 nthrds(1024, 1, 1);
895 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
896 const int nb = ((*n) + 1024 - 1)/ 1024;
897
898 hip_redbuf_check_alloc_xp(nb);
899
900 if (*n > 0) {
901 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_kernel<real, real_xp>),
902 nblcks, nthrds, 0, stream,
903 (real *) a, (real *) b,
904 (real *) c, (real_xp *) redbuf_xp.dev, *n);
905 HIP_CHECK(hipGetLastError());
906 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
907 1, 1024, 0, stream, (real_xp * ) redbuf_xp.dev, nb);
908 HIP_CHECK(hipGetLastError());
909 }
910 else {
911 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
912 }
913 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
914 (real_xp *) redbuf_xp.dev, 1, stream);
915
916 return ((real_xp *) redbuf_xp.host)[0];
917 }
918
923 void hip_glsc3_many(real_xp *h, void * w, void *v,void *mult, int *j, int *n,
924 hipStream_t stream){
925 int pow2 = 1;
926 while(pow2 < (*j)){
927 pow2 = 2*pow2;
928 }
929 const int nt = 1024/pow2;
930 const dim3 nthrds(pow2, nt, 1);
931 const dim3 nblcks(((*n)+nt - 1)/nt, 1, 1);
932 const dim3 nthrds_red(1024,1,1);
933 const dim3 nblcks_red( (*j),1,1);
934 const int nb = ((*n) + nt - 1)/nt;
935
936 hip_redbuf_check_alloc_xp((*j)*nb);
937
938 if (*n > 0) {
939 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_many_kernel<real, real_xp>),
940 nblcks, nthrds, 0, stream,
941 (const real *) w, (const real **) v,
942 (const real *)mult, (real_xp *) redbuf_xp.dev, *j, *n);
943 HIP_CHECK(hipGetLastError());
944
945 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_reduce_kernel<real_xp>),
946 nblcks_red, nthrds_red, 0, stream,
947 (real_xp *) redbuf_xp.dev, nb, *j);
948 HIP_CHECK(hipGetLastError());
949 }
950 else {
951 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
952 }
953 hip_global_reduce_add_xp(h, redbuf_xp.dev, (*j), stream);
954 }
955
960 real_xp hip_glsc2(void *a, void *b, int *n, hipStream_t stream) {
961
962 const dim3 nthrds(1024, 1, 1);
963 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
964 const int nb = ((*n) + 1024 - 1)/ 1024;
965
966
967 hip_redbuf_check_alloc_xp(nb);
968
969 if( *n > 0) {
970 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc2_kernel<real, real_xp>),
971 nblcks, nthrds, 0, stream,
972 (real *) a, (real *) b, (real_xp *) redbuf_xp.dev, *n);
973 HIP_CHECK(hipGetLastError());
974 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
975 1, 1024, 0, stream, (real_xp *) redbuf_xp.dev, nb);
976 HIP_CHECK(hipGetLastError());
977 }
978 else {
979 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
980 }
981 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
982 (real_xp *) redbuf_xp.dev, 1, stream);
983
984 return ((real_xp *) redbuf_xp.host)[0];
985 }
986
991 real_xp hip_glsubnorm2(void* a, void* b, int* n, hipStream_t stream) {
992
993 const dim3 nthrds(1024, 1, 1);
994 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
995 const int nb = ((*n) + 1024 - 1) / 1024;
996
997 hip_redbuf_check_alloc_xp(nb);
998
999 if (*n > 0) {
1000 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsubnorm2_kernel<real, real_xp>),
1001 nblcks, nthrds, 0, stream,
1002 (real *) a, (real *) b,
1003 (real_xp *) redbuf_xp.dev, *n);
1004 HIP_CHECK(hipGetLastError());
1005 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
1006 1, 1024, 0, stream, (real_xp *) redbuf_xp.dev, nb);
1007 HIP_CHECK(hipGetLastError());
1008 }
1009 else {
1010 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1011 }
1012 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
1013 (real_xp *) redbuf_xp.dev, 1, stream);
1014
1015 return ((real_xp *) redbuf_xp.host)[0];
1016 }
1017
1022 real_xp hip_glsum(void *a, int *n, hipStream_t stream) {
1023 const dim3 nthrds(1024, 1, 1);
1024 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1025 const int nb = ((*n) + 1024 - 1)/ 1024;
1026
1027 hip_redbuf_check_alloc_xp(nb);
1028 if( *n > 0) {
1029 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsum_kernel<real, real_xp>),
1030 nblcks, nthrds, 0, stream,
1031 (real *) a, (real_xp * ) redbuf_xp.dev, *n);
1032 HIP_CHECK(hipGetLastError());
1033 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
1034 1, 1024, 0, stream, (real_xp *) redbuf_xp.dev, nb);
1035 HIP_CHECK(hipGetLastError());
1036 }
1037 else {
1038 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1039 }
1040
1041 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
1042 redbuf_xp.dev, 1, stream);
1043
1044 return ((real_xp *) redbuf_xp.host)[0];
1045 }
1046
1047
1052 real hip_glmax(void *a, real *ninf, int *n, hipStream_t stream) {
1053 const dim3 nthrds(1024, 1, 1);
1054 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1055 const int nb = ((*n) + 1024 - 1)/ 1024;
1056
1057 hip_redbuf_check_alloc(nb);
1058 if( *n > 0) {
1059 hipLaunchKernelGGL(HIP_KERNEL_NAME(glmax_kernel<real>),
1060 nblcks, nthrds, 0, stream,
1061 (real *) a, *ninf, (real *) redbuf.dev, *n);
1062 HIP_CHECK(hipGetLastError());
1063 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_max_kernel<real>),
1064 1, 1024, 0, stream, (real *) redbuf.dev, *ninf, nb);
1065 HIP_CHECK(hipGetLastError());
1066 }
1067 else {
1068 int nel = (int) (redbuf.size / sizeof(real));
1069 hip_rzero(redbuf.dev, &nel, stream);
1070 }
1071
1072 hip_global_reduce_max(((real *) redbuf.host),
1073 (real *) redbuf.dev, 1, stream);
1074
1075 return ((real *) redbuf.host)[0];
1076 }
1077
1082 real hip_glmin(void *a, real *pinf, int *n, hipStream_t stream) {
1083 const dim3 nthrds(1024, 1, 1);
1084 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1085 const int nb = ((*n) + 1024 - 1)/ 1024;
1086
1087 hip_redbuf_check_alloc(nb);
1088 if( *n > 0) {
1089 hipLaunchKernelGGL(HIP_KERNEL_NAME(glmin_kernel<real>),
1090 nblcks, nthrds, 0, stream,
1091 (real *) a, *pinf, (real *) redbuf.dev, *n);
1092 HIP_CHECK(hipGetLastError());
1093 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_min_kernel<real>),
1094 1, 1024, 0, stream, (real *) redbuf.dev, *pinf, nb);
1095 HIP_CHECK(hipGetLastError());
1096 }
1097 else {
1098 int nel = (int) (redbuf.size / sizeof(real));
1099 hip_rzero(redbuf.dev, &nel, stream);
1100 }
1101
1102 hip_global_reduce_min(((real *) redbuf.host),
1103 (real *) redbuf.dev, 1, stream);
1104
1105 return ((real *) redbuf.host)[0];
1106 }
1107
1112 real hip_glamax(void *a, int *n, hipStream_t stream) {
1113 const dim3 nthrds(1024, 1, 1);
1114 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1115 const int nb = ((*n) + 1024 - 1)/ 1024;
1116 const real zero = 0.0;
1117
1118 hip_redbuf_check_alloc(nb);
1119 if( *n > 0) {
1120 hipLaunchKernelGGL(HIP_KERNEL_NAME(glamax_kernel<real>),
1121 nblcks, nthrds, 0, stream,
1122 (real *) a, (real *) redbuf.dev, *n);
1123 HIP_CHECK(hipGetLastError());
1124 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_max_kernel<real>),
1125 1, 1024, 0, stream, (real *) redbuf.dev, zero, nb);
1126 HIP_CHECK(hipGetLastError());
1127 }
1128 else {
1129 int nel = (int) (redbuf.size / sizeof(real));
1130 hip_rzero(redbuf.dev, &nel, stream);
1131 }
1132
1133 hip_global_reduce_max(((real *) redbuf.host),
1134 (real *) redbuf.dev, 1, stream);
1135
1136 return ((real *) redbuf.host)[0];
1137 }
1138
1143 void hip_absval(void *a, int *n, hipStream_t stream) {
1144
1145 const dim3 nthrds(1024, 1, 1);
1146 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1147
1148 hipLaunchKernelGGL(HIP_KERNEL_NAME(absval_kernel<real>),
1149 nblcks, nthrds, 0, stream, (real *) a, *n);
1150 HIP_CHECK(hipGetLastError());
1151
1152}
1153
1154 // ======================================================================== //
1155 // Point-wise operations.
1156
1161 void hip_pwmax_vec2(void* a, void* b, int* n, hipStream_t stream) {
1162
1163 const dim3 nthrds(1024, 1, 1);
1164 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1165
1166 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_vec2_kernel<real>),
1167 nblcks, nthrds, 0, stream, (real*)a, (real*)b, *n);
1168 HIP_CHECK(hipGetLastError());
1169}
1170
1175 void hip_pwmax_vec3(void *a, void *b, void *c, int *n, hipStream_t stream) {
1176
1177 const dim3 nthrds(1024, 1, 1);
1178 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1179
1180 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_vec3_kernel<real>),
1181 nblcks, nthrds, 0, stream,
1182 (real *)a, (real *)b, (real *)c, *n);
1183 HIP_CHECK(hipGetLastError());
1184 }
1185
1190 void hip_pwmax_sca2(void *a, real *c, int *n, hipStream_t stream) {
1191
1192 const dim3 nthrds(1024, 1, 1);
1193 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1194
1195 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_sca2_kernel<real>),
1196 nblcks, nthrds, 0, stream, (real *)a, *c, *n);
1197 HIP_CHECK(hipGetLastError());
1198 }
1199
1204 void hip_pwmax_sca3(void *a, void *b, real *c, int *n, hipStream_t stream) {
1205
1206 const dim3 nthrds(1024, 1, 1);
1207 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1208
1209 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_sca3_kernel<real>),
1210 nblcks, nthrds, 0, stream,
1211 (real *)a, (real *)b, *c, *n);
1212 HIP_CHECK(hipGetLastError());
1213 }
1214
1219 void hip_pwmin_vec2(void *a, void *b, int *n, hipStream_t stream) {
1220
1221 const dim3 nthrds(1024, 1, 1);
1222 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1223
1224 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_vec2_kernel<real>),
1225 nblcks, nthrds, 0, stream,
1226 (real *)a, (real *)b, *n);
1227 HIP_CHECK(hipGetLastError());
1228 }
1229
1234 void hip_pwmin_vec3(void *a, void *b, void *c, int *n, hipStream_t stream) {
1235
1236 const dim3 nthrds(1024, 1, 1);
1237 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1238
1239 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_vec3_kernel<real>),
1240 nblcks, nthrds, 0, stream,
1241 (real *)a, (real *)b, (real *)c, *n);
1242 HIP_CHECK(hipGetLastError());
1243 }
1244
1249 void hip_pwmin_sca2(void *a, real *c, int *n, hipStream_t stream) {
1250
1251 const dim3 nthrds(1024, 1, 1);
1252 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1253
1254 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_sca2_kernel<real>),
1255 nblcks, nthrds, 0, stream,
1256 (real *)a, *c, *n);
1257 HIP_CHECK(hipGetLastError());
1258 }
1259
1264 void hip_pwmin_sca3(void *a, void *b, real *c, int *n, hipStream_t stream) {
1265
1266 const dim3 nthrds(1024, 1, 1);
1267 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1268
1269 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_sca3_kernel<real>),
1270 nblcks, nthrds, 0, stream,
1271 (real *)a, (real *)b, *c, *n);
1272 HIP_CHECK(hipGetLastError());
1273 }
1274
1275 // ======================================================================== //
1276
1280 void hip_iadd(void *a, int *c, int *n, hipStream_t stream) {
1281
1282 const dim3 nthrds(1024, 1, 1);
1283 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1284
1285 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd_kernel<int>),
1286 nblcks, nthrds, 0, stream,
1287 (int *) a, *c, *n);
1288 HIP_CHECK(hipGetLastError());
1289 }
1290
1291} /* extern "C" */
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_compr_kernel.h:978
u
__global__ void const T *__restrict__ u
Definition ax_helm_compr_kernel.h:199
e
const int e
Definition ax_helm_compr_kernel.h:244
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_compr_kernel.h:977
j
const int j
Definition ax_helm_compr_kernel.h:245
x
__global__ void const T *__restrict__ x
Definition cdtp_kernel.h:110
w3
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
Definition cdtp_kernel.h:117
MASKED_ATOMIC_RED_BLOCK
#define MASKED_ATOMIC_RED_BLOCK
Definition math_kernel.h:198
device_config.h
real
double real
Definition device_config.h:12
real_xp
double real_xp
Definition device_config.h:16
device_mpi_op.h
DEVICE_MPI_MAX
#define DEVICE_MPI_MAX
Definition device_mpi_op.h:10
DEVICE_MPI_MIN
#define DEVICE_MPI_MIN
Definition device_mpi_op.h:11
DEVICE_MPI_SUM
#define DEVICE_MPI_SUM
Definition device_mpi_op.h:9
device_mpi_allreduce
void device_mpi_allreduce(void *buf_d, void *buf, int count, int nbytes, int op)
Definition device_mpi_reduce.c:46
device_mpi_reduce.h
device_nccl_op.h
DEVICE_NCCL_MAX
#define DEVICE_NCCL_MAX
Definition device_nccl_op.h:10
DEVICE_NCCL_MIN
#define DEVICE_NCCL_MIN
Definition device_nccl_op.h:11
DEVICE_NCCL_SUM
#define DEVICE_NCCL_SUM
Definition device_nccl_op.h:9
device_nccl_allreduce
void device_nccl_allreduce(void *sbuf_d, void *rbuf_d, int count, int nbytes, int op, void *stream)
Definition device_nccl_reduce.c:45
device_nccl_reduce.h
hip_buffer_reserve
void hip_buffer_reserve(hip_buffer_t *buf, size_t size)
Definition buffer.hip:47
HIP_BUFFER_INIT
#define HIP_BUFFER_INIT
Definition buffer.h:61
HIP_CHECK
#define HIP_CHECK(err)
Definition check.h:8
math_kernel.h
hip_vdot3
void hip_vdot3(void *dot, void *u1, void *u2, void *u3, void *v1, void *v2, void *v3, int *n, hipStream_t strm)
Definition math.hip:724
hip_global_reduce_add
void hip_global_reduce_add(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:777
hip_col3
void hip_col3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:617
hip_cwrap
void hip_cwrap(void *a, real *min_val, real *max_val, int *n, hipStream_t strm)
Definition math.hip:336
hip_addsqr2s2
void hip_addsqr2s2(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:492
hip_glmax
real hip_glmax(void *a, real *ninf, int *n, hipStream_t stream)
Definition math.hip:1052
hip_rzero
void hip_rzero(void *a, int *n, hipStream_t strm)
Definition math.hip:232
hip_power
void hip_power(void *ap, void *a, real *p, int *n, hipStream_t strm)
Definition math.hip:364
hip_cfill_mask
void hip_cfill_mask(void *a, real *c, int *size, void *mask, int *mask_size, hipStream_t strm)
Definition math.hip:216
hip_pwmax_sca3
void hip_pwmax_sca3(void *a, void *b, real *c, int *n, hipStream_t stream)
Definition math.hip:1204
hip_invcol3
void hip_invcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:587
hip_absval
void hip_absval(void *a, int *n, hipStream_t stream)
Definition math.hip:1143
hip_masked_gather_copy_aligned
void hip_masked_gather_copy_aligned(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:128
hip_vcross
void hip_vcross(void *u1, void *u2, void *u3, void *v1, void *v2, void *v3, void *w1, void *w2, void *w3, int *n, hipStream_t strm)
Definition math.hip:742
hip_addcol3
void hip_addcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:677
hip_add3s2
void hip_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n, hipStream_t strm)
Definition math.hip:508
hip_masked_scatter_copy_aligned
void hip_masked_scatter_copy_aligned(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:181
hip_add5s4
void hip_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1, real *c2, real *c3, real *c4, int *n, hipStream_t strm)
Definition math.hip:542
hip_masked_copy_aligned
void hip_masked_copy_aligned(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:95
hip_add2s2
void hip_add2s2(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:457
hip_add4
void hip_add4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm)
Definition math.hip:425
hip_redbuf_check_alloc_xp
void hip_redbuf_check_alloc_xp(int nb)
Definition math.hip:770
hip_cadd2
void hip_cadd2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:322
hip_glamax
real hip_glamax(void *a, int *n, hipStream_t stream)
Definition math.hip:1112
hip_vlsc3
real hip_vlsc3(void *u, void *v, void *w, int *n, hipStream_t stream)
Definition math.hip:863
hip_cdiv2
void hip_cdiv2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:294
hip_pwmin_sca2
void hip_pwmin_sca2(void *a, real *c, int *n, hipStream_t stream)
Definition math.hip:1249
hip_cmult2
void hip_cmult2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:266
hip_invcol1
void hip_invcol1(void *a, int *n, hipStream_t strm)
Definition math.hip:559
hip_invcol2
void hip_invcol2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:573
hip_global_reduce_max
void hip_global_reduce_max(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:820
hip_cdiv
void hip_cdiv(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:280
hip_glsubnorm2
real_xp hip_glsubnorm2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:991
hip_sub2
void hip_sub2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:647
hip_glsc2
real_xp hip_glsc2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:960
hip_add4s3
void hip_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2, real *c3, int *n, hipStream_t strm)
Definition math.hip:525
hip_col2
void hip_col2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:602
hip_glmin
real hip_glmin(void *a, real *pinf, int *n, hipStream_t stream)
Definition math.hip:1082
hip_pwmin_vec3
void hip_pwmin_vec3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:1234
hip_glsc3_many
void hip_glsc3_many(real_xp *h, void *w, void *v, void *mult, int *j, int *n, hipStream_t stream)
Definition math.hip:923
hip_addcol3s2
void hip_addcol3s2(void *a, void *b, void *c, real *s, int *n, hipStream_t strm)
Definition math.hip:708
hip_masked_gather_copy
void hip_masked_gather_copy(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:112
hip_cfill
void hip_cfill(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:378
hip_masked_atomic_reduction
void hip_masked_atomic_reduction(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:198
hip_masked_copy_0
void hip_masked_copy_0(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:78
hip_add3
void hip_add3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:410
hip_pwmax_vec3
void hip_pwmax_vec3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:1175
hip_add2
void hip_add2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:395
hip_sqrt_inplace
void hip_sqrt_inplace(void *a, int *n, hipStream_t strm)
Definition math.hip:351
hip_glsum
real_xp hip_glsum(void *a, int *n, hipStream_t stream)
Definition math.hip:1022
hip_glsc3
real_xp hip_glsc3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:892
hip_global_reduce_add_xp
void hip_global_reduce_add_xp(real_xp *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:798
redbuf
hip_buffer_t redbuf
Definition math.hip:763
hip_copy
void hip_copy(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:55
hip_subcol3
void hip_subcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:632
hip_pwmax_sca2
void hip_pwmax_sca2(void *a, real *c, int *n, hipStream_t stream)
Definition math.hip:1190
hip_iadd
void hip_iadd(void *a, int *c, int *n, hipStream_t stream)
Definition math.hip:1280
hip_pwmin_sca3
void hip_pwmin_sca3(void *a, void *b, real *c, int *n, hipStream_t stream)
Definition math.hip:1264
hip_add2s1
void hip_add2s1(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:441
hip_radd
void hip_radd(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:308
hip_redbuf_check_alloc
void hip_redbuf_check_alloc(int nb)
Definition math.hip:766
hip_add2s2_many
void hip_add2s2_many(void *x, void **p, void *alpha, int *j, int *n, hipStream_t strm)
Definition math.hip:474
redbuf_xp
hip_buffer_t redbuf_xp
Definition math.hip:764
hip_addcol4
void hip_addcol4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm)
Definition math.hip:692
hip_pwmax_vec2
void hip_pwmax_vec2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:1161
hip_pwmin_vec2
void hip_pwmin_vec2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:1219
hip_sub3
void hip_sub3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:662
hip_face_masked_gather_copy
void hip_face_masked_gather_copy(void *a, void *b, void *mask, void *facet, int *n1, int *n2, int *lx, int *ly, int *lz, int *m, hipStream_t strm)
Definition math.hip:144
hip_masked_scatter_copy
void hip_masked_scatter_copy(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:164
hip_cmult
void hip_cmult(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:252
hip_global_reduce_min
void hip_global_reduce_min(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:841
mask
Object for handling masks in Neko.
Definition mask.f90:34
hip_buffer
Definition buffer.h:51
hip_buffer::host
void * host
Definition buffer.h:52
hip_buffer::size
size_t size
Definition buffer.h:54
hip_buffer::dev
void * dev
Definition buffer.h:53
hip_zerocopy
int hip_zerocopy(void)
Definition unified.hip:122