Neko 1.99.6
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-2025, 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(1024, 1, 1);
202 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
203
204 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_atomic_reduction_kernel<real>),
205 nblcks, nthrds, 0, strm, (real *) a,
206 (real *) b, (int *) mask, *n, *m);
207
208 HIP_CHECK(hipGetLastError());
209
210 }
211
215 void hip_cfill_mask(void* a, real* c, int* size, void* mask, int* mask_size,
216 hipStream_t strm) {
217
218 const dim3 nthrds(1024, 1, 1);
219 const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
220
221 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_mask_kernel<real>),
222 nblcks, nthrds, 0, strm, (real*)a,
223 *c, *size, (int*)mask, *mask_size);
224
225 HIP_CHECK(hipGetLastError());
226 }
227
231 void hip_rzero(void *a, int *n, hipStream_t strm) {
232 if (hip_zerocopy()) {
233 if (*n == 0) return;
234 /* a may alias pageable host memory, which hipMemset rejects;
235 zero with a kernel instead (works on any pointer) */
236 const dim3 nthrds(1024, 1, 1);
237 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
238
239 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_kernel<real>),
240 nblcks, nthrds, 0, strm,
241 (real *) a, (real) 0.0, *n);
242 HIP_CHECK(hipGetLastError());
243 return;
244 }
245 HIP_CHECK(hipMemsetAsync(a, 0, (*n) * sizeof(real), strm));
246 }
247
251 void hip_cmult(void *a, real *c, int *n, hipStream_t strm) {
252
253 const dim3 nthrds(1024, 1, 1);
254 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
255
256 hipLaunchKernelGGL(HIP_KERNEL_NAME(cmult_kernel<real>),
257 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
258 HIP_CHECK(hipGetLastError());
259
260 }
261
265 void hip_cmult2(void *a, void *b, real *c, int *n, hipStream_t strm) {
266
267 const dim3 nthrds(1024, 1, 1);
268 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
269
270 hipLaunchKernelGGL(HIP_KERNEL_NAME(cmult2_kernel<real>),
271 nblcks, nthrds, 0, strm, (real *) a,(real *) b, *c, *n);
272 HIP_CHECK(hipGetLastError());
273
274 }
275
279 void hip_cdiv(void *a, real *c, int *n, hipStream_t strm) {
280
281 const dim3 nthrds(1024, 1, 1);
282 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
283
284 hipLaunchKernelGGL(HIP_KERNEL_NAME(cdiv_kernel<real>),
285 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
286 HIP_CHECK(hipGetLastError());
287
288 }
289
293 void hip_cdiv2(void *a, void *b, real *c, int *n, hipStream_t strm) {
294
295 const dim3 nthrds(1024, 1, 1);
296 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
297
298 hipLaunchKernelGGL(HIP_KERNEL_NAME(cdiv2_kernel<real>),
299 nblcks, nthrds, 0, strm, (real *) a,(real *) b, *c, *n);
300 HIP_CHECK(hipGetLastError());
301
302 }
303
307 void hip_radd(void *a, real *c, int *n, hipStream_t strm) {
308
309 const dim3 nthrds(1024, 1, 1);
310 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
311
312 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd_kernel<real>),
313 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
314 HIP_CHECK(hipGetLastError());
315 }
316
321 void hip_cadd2(void *a, void *b, real *c, int *n, hipStream_t strm) {
322
323 const dim3 nthrds(1024, 1, 1);
324 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
325
326 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd2_kernel<real>),
327 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *c, *n);
328 HIP_CHECK(hipGetLastError());
329 }
330
335 void hip_cwrap(void *a, real *min_val, real *max_val, int *n,
336 hipStream_t strm) {
337
338 const dim3 nthrds(1024, 1, 1);
339 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
340
341 hipLaunchKernelGGL(HIP_KERNEL_NAME(cwrap_kernel<real>),
342 nblcks, nthrds, 0, strm, (real *) a,
343 *min_val, *max_val, *n);
344 HIP_CHECK(hipGetLastError());
345 }
346
350 void hip_sqrt_inplace(void *a, int *n, hipStream_t strm) {
351
352 const dim3 nthrds(1024, 1, 1);
353 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
354
355 hipLaunchKernelGGL(HIP_KERNEL_NAME(sqrt_inplace_kernel<real>),
356 nblcks, nthrds, 0, strm, (real *) a, *n);
357 HIP_CHECK(hipGetLastError());
358 }
359
363 void hip_power(void *ap, void *a, real *p, int *n, hipStream_t strm) {
364
365 const dim3 nthrds(1024, 1, 1);
366 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
367
368 hipLaunchKernelGGL(HIP_KERNEL_NAME(power_kernel<real>),
369 nblcks, nthrds, 0, strm,
370 (real *) ap, (real *) a, *p, *n);
371 HIP_CHECK(hipGetLastError());
372 }
373
377 void hip_cfill(void *a, real *c, int *n, hipStream_t strm) {
378
379 const dim3 nthrds(1024, 1, 1);
380 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
381
382 if (*n > 0) {
383 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_kernel<real>),
384 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
385 HIP_CHECK(hipGetLastError());
386 }
387
388 }
389
394 void hip_add2(void *a, void *b, int *n, hipStream_t strm) {
395
396 const dim3 nthrds(1024, 1, 1);
397 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
398
399 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2_kernel<real>),
400 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *n);
401 HIP_CHECK(hipGetLastError());
402
403 }
404
409 void hip_add3(void *a, void *b, void *c, int *n, hipStream_t strm) {
410
411 const dim3 nthrds(1024, 1, 1);
412 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
413
414 hipLaunchKernelGGL(HIP_KERNEL_NAME(add3_kernel<real>),
415 nblcks, nthrds, 0, strm, (real *) a,
416 (real *) b, (real *) c, *n);
417 HIP_CHECK(hipGetLastError());
418 }
419
424 void hip_add4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm) {
425
426 const dim3 nthrds(1024, 1, 1);
427 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
428
429 hipLaunchKernelGGL(HIP_KERNEL_NAME(add4_kernel<real>),
430 nblcks, nthrds, 0, strm, (real *) a,
431 (real *) b, (real *) c, (real *) d, *n);
432 HIP_CHECK(hipGetLastError());
433 }
434
440 void hip_add2s1(void *a, void *b, real *c1, int *n, hipStream_t strm) {
441
442 const dim3 nthrds(1024, 1, 1);
443 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
444
445 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s1_kernel<real>),
446 nblcks, nthrds, 0, strm, (real *) a,
447 (real *) b, *c1, *n);
448 HIP_CHECK(hipGetLastError());
449 }
450
456 void hip_add2s2(void *a, void *b, real *c1, int *n, hipStream_t strm) {
457
458 const dim3 nthrds(1024, 1, 1);
459 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
460
461 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s2_kernel<real>),
462 nblcks, nthrds, 0, strm, (real *) a, (real *) b,
463 *c1, *n);
464 HIP_CHECK(hipGetLastError());
465 }
466
473 void hip_add2s2_many(void *x, void **p, void *alpha, int *j, int *n,
474 hipStream_t strm) {
475
476 const dim3 nthrds(1024, 1, 1);
477 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
478
479 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s2_many_kernel<real>),
480 nblcks, nthrds, 0, strm, (real *) x,
481 (const real **) p, (real *) alpha, *j, *n);
482 HIP_CHECK(hipGetLastError());
483
484 }
485
491 void hip_addsqr2s2(void *a, void *b, real *c1, int *n, hipStream_t strm) {
492
493 const dim3 nthrds(1024, 1, 1);
494 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
495
496 hipLaunchKernelGGL(HIP_KERNEL_NAME(addsqr2s2_kernel<real>),
497 nblcks, nthrds, 0, strm, (real *) a,
498 (real *) b, *c1, *n);
499 HIP_CHECK(hipGetLastError());
500 }
501
507 void hip_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n,
508 hipStream_t strm) {
509
510 const dim3 nthrds(1024, 1, 1);
511 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
512
513 hipLaunchKernelGGL(HIP_KERNEL_NAME(add3s2_kernel<real>),
514 nblcks, nthrds, 0, strm, (real *) a,
515 (real *) b, (real *) c, *c1, *c2, *n);
516 HIP_CHECK(hipGetLastError());
517 }
518
524 void hip_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2,
525 real *c3, int *n, hipStream_t strm) {
526
527 const dim3 nthrds(1024, 1, 1);
528 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
529
530 hipLaunchKernelGGL(HIP_KERNEL_NAME(add4s3_kernel<real>),
531 nblcks, nthrds, 0, strm, (real *) a,
532 (real *) b, (real *) c, (real *) d, *c1, *c2, *c3, *n);
533 HIP_CHECK(hipGetLastError());
534 }
535
541 void hip_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1,
542 real *c2, real *c3, real *c4, int *n, hipStream_t strm) {
543
544 const dim3 nthrds(1024, 1, 1);
545 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
546
547 hipLaunchKernelGGL(HIP_KERNEL_NAME(add5s4_kernel<real>),
548 nblcks, nthrds, 0, strm, (real *) a,
549 (real *) b, (real *) c, (real *) d, (real *) e,
550 *c1, *c2, *c3, *c4, *n);
551 HIP_CHECK(hipGetLastError());
552 }
553
558 void hip_invcol1(void *a, int *n, hipStream_t strm) {
559
560 const dim3 nthrds(1024, 1, 1);
561 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
562
563 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol1_kernel<real>),
564 nblcks, nthrds, 0, strm, (real *) a, *n);
565 HIP_CHECK(hipGetLastError());
566 }
567
572 void hip_invcol2(void *a, void *b, int *n, hipStream_t strm) {
573
574 const dim3 nthrds(1024, 1, 1);
575 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
576
577 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol2_kernel<real>),
578 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *n);
579 HIP_CHECK(hipGetLastError());
580 }
581
586 void hip_invcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
587
588 const dim3 nthrds(1024, 1, 1);
589 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
590
591 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol3_kernel<real>),
592 nblcks, nthrds, 0, strm, (real *) a,
593 (real *) b, (real *) c, *n);
594 HIP_CHECK(hipGetLastError());
595 }
596
601 void hip_col2(void *a, void *b, int *n, hipStream_t strm) {
602
603 const dim3 nthrds(1024, 1, 1);
604 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
605
606 hipLaunchKernelGGL(HIP_KERNEL_NAME(col2_kernel<real>),
607 nblcks, nthrds, 0, strm, (real *) a,
608 (real *) b, *n);
609 HIP_CHECK(hipGetLastError());
610 }
611
616 void hip_col3(void *a, void *b, void *c, int *n, hipStream_t strm) {
617
618 const dim3 nthrds(1024, 1, 1);
619 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
620
621 hipLaunchKernelGGL(HIP_KERNEL_NAME(col3_kernel<real>),
622 nblcks, nthrds, 0, strm, (real *) a,
623 (real *) b, (real *) c, *n);
624 HIP_CHECK(hipGetLastError());
625 }
626
631 void hip_subcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
632
633 const dim3 nthrds(1024, 1, 1);
634 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
635
636 hipLaunchKernelGGL(HIP_KERNEL_NAME(subcol3_kernel<real>),
637 nblcks, nthrds, 0, strm, (real *) a,
638 (real *) b, (real *) c, *n);
639 HIP_CHECK(hipGetLastError());
640 }
641
646 void hip_sub2(void *a, void *b, int *n, hipStream_t strm) {
647
648 const dim3 nthrds(1024, 1, 1);
649 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
650
651 hipLaunchKernelGGL(HIP_KERNEL_NAME(sub2_kernel<real>),
652 nblcks, nthrds, 0, strm, (real *) a,
653 (real *) b, *n);
654 HIP_CHECK(hipGetLastError());
655 }
656
661 void hip_sub3(void *a, void *b, void *c, int *n, hipStream_t strm) {
662
663 const dim3 nthrds(1024, 1, 1);
664 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
665
666 hipLaunchKernelGGL(HIP_KERNEL_NAME(sub3_kernel<real>),
667 nblcks, nthrds, 0, strm, (real *) a,
668 (real *) b, (real *) c, *n);
669 HIP_CHECK(hipGetLastError());
670 }
671
676 void hip_addcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
677
678 const dim3 nthrds(1024, 1, 1);
679 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
680
681 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol3_kernel<real>),
682 nblcks, nthrds, 0, strm, (real *) a,
683 (real *) b, (real *) c, *n);
684 HIP_CHECK(hipGetLastError());
685 }
686
691 void hip_addcol4(void *a, void *b, void *c, void *d, int *n,
692 hipStream_t strm) {
693
694 const dim3 nthrds(1024, 1, 1);
695 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
696
697 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol4_kernel<real>),
698 nblcks, nthrds, 0, strm, (real *) a,
699 (real *) b, (real *) c, (real *) d, *n);
700 HIP_CHECK(hipGetLastError());
701 }
702
707 void hip_addcol3s2(void *a, void *b, void *c, real *s, int *n,
708 hipStream_t strm) {
709
710 const dim3 nthrds(1024, 1, 1);
711 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
712
713 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol3s2_kernel<real>),
714 nblcks, nthrds, 0, strm, (real *) a,
715 (real *) b, (real *) c, *s, *n);
716 HIP_CHECK(hipGetLastError());
717 }
718
723 void hip_vdot3(void *dot, void *u1, void *u2, void *u3,
724 void *v1, void *v2, void *v3, int *n,
725 hipStream_t strm) {
726
727 const dim3 nthrds(1024, 1, 1);
728 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
729
730 hipLaunchKernelGGL(HIP_KERNEL_NAME(vdot3_kernel<real>),
731 nblcks, nthrds, 0, strm,
732 (real *) dot, (real *) u1, (real *) u2, (real *) u3,
733 (real *) v1, (real *) v2, (real *) v3, *n);
734 HIP_CHECK(hipGetLastError());
735 }
736
741 void hip_vcross(void *u1, void *u2, void *u3,
742 void *v1, void *v2, void *v3,
743 void *w1, void *w2, void *w3,
744 int *n, hipStream_t strm) {
745
746 const dim3 nthrds(1024, 1, 1);
747 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
748
749 hipLaunchKernelGGL(HIP_KERNEL_NAME(vcross_kernel<real>),
750 nblcks, nthrds, 0, strm,
751 (real *) u1, (real *) u2, (real *) u3,
752 (real *) v1, (real *) v2, (real *) v3,
753 (real *) w1, (real *) w2, (real *) w3, *n);
754 HIP_CHECK(hipGetLastError());
755 }
756
757
758 /*
759 * Reduction buffers, owned by the device layer and released
760 * on device teardown (hip_buffer_free_all in hip_finalize)
761 */
762 hip_buffer_t redbuf = HIP_BUFFER_INIT;
763 hip_buffer_t redbuf_xp = HIP_BUFFER_INIT;
764
765 void hip_redbuf_check_alloc(int nb) {
766 hip_buffer_reserve(&redbuf, (nb + 1) * sizeof(real));
767 }
768
769 void hip_redbuf_check_alloc_xp(int nb) {
770 hip_buffer_reserve(&redbuf_xp, (nb + 1) * sizeof(real_xp));
771 }
772
776 void hip_global_reduce_add(real * bufred, void * bufred_d, int n,
777 const hipStream_t stream) {
778 #ifdef HAVE_RCCL
779 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
780 DEVICE_NCCL_SUM, stream);
781 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
782 hipMemcpyDeviceToHost, stream));
783 HIP_CHECK(hipStreamSynchronize(stream));
784 #elif HAVE_DEVICE_MPI
785 HIP_CHECK(hipStreamSynchronize(stream));
786 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_SUM);
787 #else
788 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
789 hipMemcpyDeviceToHost, stream));
790 HIP_CHECK(hipStreamSynchronize(stream));
791 #endif
792 }
793
797 void hip_global_reduce_add_xp(real_xp * bufred, void * bufred_d, int n,
798 const hipStream_t stream) {
799 #ifdef HAVE_RCCL
800 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real_xp),
801 DEVICE_NCCL_SUM, stream);
802 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real_xp),
803 hipMemcpyDeviceToHost, stream));
804 HIP_CHECK(hipStreamSynchronize(stream));
805 #elif HAVE_DEVICE_MPI
806 HIP_CHECK(hipStreamSynchronize(stream));
807 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real_xp), DEVICE_MPI_SUM);
808 #else
809 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real_xp),
810 hipMemcpyDeviceToHost, stream));
811 HIP_CHECK(hipStreamSynchronize(stream));
812 #endif
813 }
814
815
819 void hip_global_reduce_max(real * bufred, void * bufred_d, int n,
820 const hipStream_t stream) {
821 #ifdef HAVE_RCCL
822 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
823 DEVICE_NCCL_MAX, stream);
824 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
825 hipMemcpyDeviceToHost, stream));
826 HIP_CHECK(hipStreamSynchronize(stream));
827 #elif HAVE_DEVICE_MPI
828 HIP_CHECK(hipStreamSynchronize(stream));
829 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MAX);
830 #else
831 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
832 hipMemcpyDeviceToHost, stream));
833 HIP_CHECK(hipStreamSynchronize(stream));
834 #endif
835 }
836
840 void hip_global_reduce_min(real * bufred, void * bufred_d, int n,
841 const hipStream_t stream) {
842 #ifdef HAVE_RCCL
843 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
844 DEVICE_NCCL_MIN, stream);
845 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
846 hipMemcpyDeviceToHost, stream));
847 HIP_CHECK(hipStreamSynchronize(stream));
848 #elif HAVE_DEVICE_MPI
849 HIP_CHECK(hipStreamSynchronize(stream));
850 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MIN);
851 #else
852 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
853 hipMemcpyDeviceToHost, stream));
854 HIP_CHECK(hipStreamSynchronize(stream));
855 #endif
856 }
857
862 real hip_vlsc3(void *u, void *v, void *w, int *n, hipStream_t stream) {
863
864 const dim3 nthrds(1024, 1, 1);
865 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
866 const int nb = ((*n) + 1024 - 1)/ 1024;
867
868 hip_redbuf_check_alloc(nb);
869
870
871 hipLaunchKernelGGL(HIP_KERNEL_NAME(vlsc3_kernel<real>),
872 nblcks, nthrds, 0, stream,
873 (real *) u, (real *) v,
874 (real *) w, (real *) redbuf.dev, *n);
875 HIP_CHECK(hipGetLastError());
876 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
877 1, 1024, 0, stream, (real *) redbuf.dev, nb);
878 HIP_CHECK(hipGetLastError());
879
880 HIP_CHECK(hipMemcpyAsync(((real *) redbuf.host), (real *) redbuf.dev,
881 sizeof(real), hipMemcpyDeviceToHost, stream));
882 HIP_CHECK(hipStreamSynchronize(stream));
883
884 return ((real *) redbuf.host)[0];
885 }
886
891 real_xp hip_glsc3(void *a, void *b, void *c, int *n, hipStream_t stream) {
892
893 const dim3 nthrds(1024, 1, 1);
894 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
895 const int nb = ((*n) + 1024 - 1)/ 1024;
896
897 hip_redbuf_check_alloc_xp(nb);
898
899 if (*n > 0) {
900 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_kernel<real, real_xp>),
901 nblcks, nthrds, 0, stream,
902 (real *) a, (real *) b,
903 (real *) c, (real_xp *) redbuf_xp.dev, *n);
904 HIP_CHECK(hipGetLastError());
905 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
906 1, 1024, 0, stream, (real_xp * ) redbuf_xp.dev, nb);
907 HIP_CHECK(hipGetLastError());
908 }
909 else {
910 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
911 }
912 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
913 (real_xp *) redbuf_xp.dev, 1, stream);
914
915 return ((real_xp *) redbuf_xp.host)[0];
916 }
917
922 void hip_glsc3_many(real_xp *h, void * w, void *v,void *mult, int *j, int *n,
923 hipStream_t stream){
924 int pow2 = 1;
925 while(pow2 < (*j)){
926 pow2 = 2*pow2;
927 }
928 const int nt = 1024/pow2;
929 const dim3 nthrds(pow2, nt, 1);
930 const dim3 nblcks(((*n)+nt - 1)/nt, 1, 1);
931 const dim3 nthrds_red(1024,1,1);
932 const dim3 nblcks_red( (*j),1,1);
933 const int nb = ((*n) + nt - 1)/nt;
934
935 hip_redbuf_check_alloc_xp((*j)*nb);
936
937 if (*n > 0) {
938 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_many_kernel<real, real_xp>),
939 nblcks, nthrds, 0, stream,
940 (const real *) w, (const real **) v,
941 (const real *)mult, (real_xp *) redbuf_xp.dev, *j, *n);
942 HIP_CHECK(hipGetLastError());
943
944 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_reduce_kernel<real_xp>),
945 nblcks_red, nthrds_red, 0, stream,
946 (real_xp *) redbuf_xp.dev, nb, *j);
947 HIP_CHECK(hipGetLastError());
948 }
949 else {
950 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
951 }
952 hip_global_reduce_add_xp(h, redbuf_xp.dev, (*j), stream);
953 }
954
959 real_xp hip_glsc2(void *a, void *b, int *n, hipStream_t stream) {
960
961 const dim3 nthrds(1024, 1, 1);
962 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
963 const int nb = ((*n) + 1024 - 1)/ 1024;
964
965
966 hip_redbuf_check_alloc_xp(nb);
967
968 if( *n > 0) {
969 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc2_kernel<real, real_xp>),
970 nblcks, nthrds, 0, stream,
971 (real *) a, (real *) b, (real_xp *) redbuf_xp.dev, *n);
972 HIP_CHECK(hipGetLastError());
973 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
974 1, 1024, 0, stream, (real_xp *) redbuf_xp.dev, nb);
975 HIP_CHECK(hipGetLastError());
976 }
977 else {
978 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
979 }
980 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
981 (real_xp *) redbuf_xp.dev, 1, stream);
982
983 return ((real_xp *) redbuf_xp.host)[0];
984 }
985
990 real_xp hip_glsubnorm2(void* a, void* b, int* n, hipStream_t stream) {
991
992 const dim3 nthrds(1024, 1, 1);
993 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
994 const int nb = ((*n) + 1024 - 1) / 1024;
995
996 hip_redbuf_check_alloc_xp(nb);
997
998 if (*n > 0) {
999 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsubnorm2_kernel<real, real_xp>),
1000 nblcks, nthrds, 0, stream,
1001 (real *) a, (real *) b,
1002 (real_xp *) redbuf_xp.dev, *n);
1003 HIP_CHECK(hipGetLastError());
1004 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
1005 1, 1024, 0, stream, (real_xp *) redbuf_xp.dev, nb);
1006 HIP_CHECK(hipGetLastError());
1007 }
1008 else {
1009 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1010 }
1011 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
1012 (real_xp *) redbuf_xp.dev, 1, stream);
1013
1014 return ((real_xp *) redbuf_xp.host)[0];
1015 }
1016
1021 real_xp hip_glsum(void *a, int *n, hipStream_t stream) {
1022 const dim3 nthrds(1024, 1, 1);
1023 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1024 const int nb = ((*n) + 1024 - 1)/ 1024;
1025
1026 hip_redbuf_check_alloc_xp(nb);
1027 if( *n > 0) {
1028 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsum_kernel<real, real_xp>),
1029 nblcks, nthrds, 0, stream,
1030 (real *) a, (real_xp * ) redbuf_xp.dev, *n);
1031 HIP_CHECK(hipGetLastError());
1032 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real_xp>),
1033 1, 1024, 0, stream, (real_xp *) redbuf_xp.dev, nb);
1034 HIP_CHECK(hipGetLastError());
1035 }
1036 else {
1037 HIP_CHECK(hipMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1038 }
1039
1040 hip_global_reduce_add_xp(((real_xp *) redbuf_xp.host),
1041 redbuf_xp.dev, 1, stream);
1042
1043 return ((real_xp *) redbuf_xp.host)[0];
1044 }
1045
1046
1051 real hip_glmax(void *a, real *ninf, int *n, hipStream_t stream) {
1052 const dim3 nthrds(1024, 1, 1);
1053 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1054 const int nb = ((*n) + 1024 - 1)/ 1024;
1055
1056 hip_redbuf_check_alloc(nb);
1057 if( *n > 0) {
1058 hipLaunchKernelGGL(HIP_KERNEL_NAME(glmax_kernel<real>),
1059 nblcks, nthrds, 0, stream,
1060 (real *) a, *ninf, (real *) redbuf.dev, *n);
1061 HIP_CHECK(hipGetLastError());
1062 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_max_kernel<real>),
1063 1, 1024, 0, stream, (real *) redbuf.dev, *ninf, nb);
1064 HIP_CHECK(hipGetLastError());
1065 }
1066 else {
1067 int nel = (int) (redbuf.size / sizeof(real));
1068 hip_rzero(redbuf.dev, &nel, stream);
1069 }
1070
1071 hip_global_reduce_max(((real *) redbuf.host),
1072 (real *) redbuf.dev, 1, stream);
1073
1074 return ((real *) redbuf.host)[0];
1075 }
1076
1081 real hip_glmin(void *a, real *pinf, int *n, hipStream_t stream) {
1082 const dim3 nthrds(1024, 1, 1);
1083 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1084 const int nb = ((*n) + 1024 - 1)/ 1024;
1085
1086 hip_redbuf_check_alloc(nb);
1087 if( *n > 0) {
1088 hipLaunchKernelGGL(HIP_KERNEL_NAME(glmin_kernel<real>),
1089 nblcks, nthrds, 0, stream,
1090 (real *) a, *pinf, (real *) redbuf.dev, *n);
1091 HIP_CHECK(hipGetLastError());
1092 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_min_kernel<real>),
1093 1, 1024, 0, stream, (real *) redbuf.dev, *pinf, nb);
1094 HIP_CHECK(hipGetLastError());
1095 }
1096 else {
1097 int nel = (int) (redbuf.size / sizeof(real));
1098 hip_rzero(redbuf.dev, &nel, stream);
1099 }
1100
1101 hip_global_reduce_min(((real *) redbuf.host),
1102 (real *) redbuf.dev, 1, stream);
1103
1104 return ((real *) redbuf.host)[0];
1105 }
1106
1111 void hip_absval(void *a, int *n, hipStream_t stream) {
1112
1113 const dim3 nthrds(1024, 1, 1);
1114 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1115
1116 hipLaunchKernelGGL(HIP_KERNEL_NAME(absval_kernel<real>),
1117 nblcks, nthrds, 0, stream, (real *) a, *n);
1118 HIP_CHECK(hipGetLastError());
1119
1120}
1121
1122 // ======================================================================== //
1123 // Point-wise operations.
1124
1129 void hip_pwmax_vec2(void* a, void* b, int* n, hipStream_t stream) {
1130
1131 const dim3 nthrds(1024, 1, 1);
1132 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1133
1134 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_vec2_kernel<real>),
1135 nblcks, nthrds, 0, stream, (real*)a, (real*)b, *n);
1136 HIP_CHECK(hipGetLastError());
1137}
1138
1143 void hip_pwmax_vec3(void *a, void *b, void *c, 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(pwmax_vec3_kernel<real>),
1149 nblcks, nthrds, 0, stream,
1150 (real *)a, (real *)b, (real *)c, *n);
1151 HIP_CHECK(hipGetLastError());
1152 }
1153
1158 void hip_pwmax_sca2(void *a, real *c, int *n, hipStream_t stream) {
1159
1160 const dim3 nthrds(1024, 1, 1);
1161 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1162
1163 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_sca2_kernel<real>),
1164 nblcks, nthrds, 0, stream, (real *)a, *c, *n);
1165 HIP_CHECK(hipGetLastError());
1166 }
1167
1172 void hip_pwmax_sca3(void *a, void *b, real *c, int *n, hipStream_t stream) {
1173
1174 const dim3 nthrds(1024, 1, 1);
1175 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1176
1177 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_sca3_kernel<real>),
1178 nblcks, nthrds, 0, stream,
1179 (real *)a, (real *)b, *c, *n);
1180 HIP_CHECK(hipGetLastError());
1181 }
1182
1187 void hip_pwmin_vec2(void *a, void *b, int *n, hipStream_t stream) {
1188
1189 const dim3 nthrds(1024, 1, 1);
1190 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1191
1192 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_vec2_kernel<real>),
1193 nblcks, nthrds, 0, stream,
1194 (real *)a, (real *)b, *n);
1195 HIP_CHECK(hipGetLastError());
1196 }
1197
1202 void hip_pwmin_vec3(void *a, void *b, void *c, int *n, hipStream_t stream) {
1203
1204 const dim3 nthrds(1024, 1, 1);
1205 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1206
1207 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_vec3_kernel<real>),
1208 nblcks, nthrds, 0, stream,
1209 (real *)a, (real *)b, (real *)c, *n);
1210 HIP_CHECK(hipGetLastError());
1211 }
1212
1217 void hip_pwmin_sca2(void *a, real *c, int *n, hipStream_t stream) {
1218
1219 const dim3 nthrds(1024, 1, 1);
1220 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1221
1222 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_sca2_kernel<real>),
1223 nblcks, nthrds, 0, stream,
1224 (real *)a, *c, *n);
1225 HIP_CHECK(hipGetLastError());
1226 }
1227
1232 void hip_pwmin_sca3(void *a, void *b, real *c, int *n, hipStream_t stream) {
1233
1234 const dim3 nthrds(1024, 1, 1);
1235 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1236
1237 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_sca3_kernel<real>),
1238 nblcks, nthrds, 0, stream,
1239 (real *)a, (real *)b, *c, *n);
1240 HIP_CHECK(hipGetLastError());
1241 }
1242
1243 // ======================================================================== //
1244
1248 void hip_iadd(void *a, int *c, int *n, hipStream_t stream) {
1249
1250 const dim3 nthrds(1024, 1, 1);
1251 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1252
1253 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd_kernel<int>),
1254 nblcks, nthrds, 0, stream,
1255 (int *) a, *c, *n);
1256 HIP_CHECK(hipGetLastError());
1257 }
1258
1259} /* 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_full_kernel.h:48
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
e
const int e
Definition ax_helm_full_kernel.h:93
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
j
const int j
Definition ax_helm_full_kernel.h:94
x
__global__ void const T *__restrict__ x
Definition cdtp_kernel.h:106
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:113
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:723
hip_global_reduce_add
void hip_global_reduce_add(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:776
hip_col3
void hip_col3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:616
hip_cwrap
void hip_cwrap(void *a, real *min_val, real *max_val, int *n, hipStream_t strm)
Definition math.hip:335
hip_addsqr2s2
void hip_addsqr2s2(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:491
hip_glmax
real hip_glmax(void *a, real *ninf, int *n, hipStream_t stream)
Definition math.hip:1051
hip_rzero
void hip_rzero(void *a, int *n, hipStream_t strm)
Definition math.hip:231
hip_power
void hip_power(void *ap, void *a, real *p, int *n, hipStream_t strm)
Definition math.hip:363
hip_cfill_mask
void hip_cfill_mask(void *a, real *c, int *size, void *mask, int *mask_size, hipStream_t strm)
Definition math.hip:215
hip_pwmax_sca3
void hip_pwmax_sca3(void *a, void *b, real *c, int *n, hipStream_t stream)
Definition math.hip:1172
hip_invcol3
void hip_invcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:586
hip_absval
void hip_absval(void *a, int *n, hipStream_t stream)
Definition math.hip:1111
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:741
hip_addcol3
void hip_addcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:676
hip_add3s2
void hip_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n, hipStream_t strm)
Definition math.hip:507
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:541
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:456
hip_add4
void hip_add4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm)
Definition math.hip:424
hip_redbuf_check_alloc_xp
void hip_redbuf_check_alloc_xp(int nb)
Definition math.hip:769
hip_cadd2
void hip_cadd2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:321
hip_vlsc3
real hip_vlsc3(void *u, void *v, void *w, int *n, hipStream_t stream)
Definition math.hip:862
hip_cdiv2
void hip_cdiv2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:293
hip_pwmin_sca2
void hip_pwmin_sca2(void *a, real *c, int *n, hipStream_t stream)
Definition math.hip:1217
hip_cmult2
void hip_cmult2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:265
hip_invcol1
void hip_invcol1(void *a, int *n, hipStream_t strm)
Definition math.hip:558
hip_invcol2
void hip_invcol2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:572
hip_global_reduce_max
void hip_global_reduce_max(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:819
hip_cdiv
void hip_cdiv(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:279
hip_glsubnorm2
real_xp hip_glsubnorm2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:990
hip_sub2
void hip_sub2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:646
hip_glsc2
real_xp hip_glsc2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:959
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:524
hip_col2
void hip_col2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:601
hip_glmin
real hip_glmin(void *a, real *pinf, int *n, hipStream_t stream)
Definition math.hip:1081
hip_pwmin_vec3
void hip_pwmin_vec3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:1202
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:922
hip_addcol3s2
void hip_addcol3s2(void *a, void *b, void *c, real *s, int *n, hipStream_t strm)
Definition math.hip:707
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:377
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:409
hip_pwmax_vec3
void hip_pwmax_vec3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:1143
hip_add2
void hip_add2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:394
hip_sqrt_inplace
void hip_sqrt_inplace(void *a, int *n, hipStream_t strm)
Definition math.hip:350
hip_glsum
real_xp hip_glsum(void *a, int *n, hipStream_t stream)
Definition math.hip:1021
hip_glsc3
real_xp hip_glsc3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:891
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:797
redbuf
hip_buffer_t redbuf
Definition math.hip:762
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:631
hip_pwmax_sca2
void hip_pwmax_sca2(void *a, real *c, int *n, hipStream_t stream)
Definition math.hip:1158
hip_iadd
void hip_iadd(void *a, int *c, int *n, hipStream_t stream)
Definition math.hip:1248
hip_pwmin_sca3
void hip_pwmin_sca3(void *a, void *b, real *c, int *n, hipStream_t stream)
Definition math.hip:1232
hip_add2s1
void hip_add2s1(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:440
hip_radd
void hip_radd(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:307
hip_redbuf_check_alloc
void hip_redbuf_check_alloc(int nb)
Definition math.hip:765
hip_add2s2_many
void hip_add2s2_many(void *x, void **p, void *alpha, int *j, int *n, hipStream_t strm)
Definition math.hip:473
redbuf_xp
hip_buffer_t redbuf_xp
Definition math.hip:763
hip_addcol4
void hip_addcol4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm)
Definition math.hip:691
hip_pwmax_vec2
void hip_pwmax_vec2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:1129
hip_pwmin_vec2
void hip_pwmin_vec2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:1187
hip_sub3
void hip_sub3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:661
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:251
hip_global_reduce_min
void hip_global_reduce_min(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:840
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