Neko 1.99.1
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 "math_kernel.h"
39
40extern "C" {
41
42#include <math/bcknd/device/device_mpi_reduce.h>
43#include <math/bcknd/device/device_mpi_op.h>
44
45#ifdef HAVE_RCCL
46#include <math/bcknd/device/device_nccl_reduce.h>
47#include <math/bcknd/device/device_nccl_op.h>
48#endif
49
53 void hip_copy(void *a, void *b, int *n, hipStream_t strm) {
54 HIP_CHECK(hipMemcpyAsync(a, b, (*n) * sizeof(real),
55 hipMemcpyDeviceToDevice, strm));
56 }
57
61 void hip_masked_copy(void *a, void *b, void *mask,
62 int *n, int *m, hipStream_t strm) {
63
64 const dim3 nthrds(1024, 1, 1);
65 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
66
67 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_copy_kernel<real>),
68 nblcks, nthrds, 0, strm, (real *) a,
69 (real *) b, (int *) mask, *n, *m);
70
71 HIP_CHECK(hipGetLastError());
72
73 }
74
78 void hip_masked_gather_copy(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_gather_copy_kernel<real>),
85 nblcks, nthrds, 0, strm, (real *) a,
86 (real *) b, (int *) mask, *n, *m);
87 HIP_CHECK(hipGetLastError());
88
89 }
90
94 void hip_masked_scatter_copy(void *a, void *b, void *mask,
95 int *n, int *m, hipStream_t strm) {
96
97 const dim3 nthrds(1024, 1, 1);
98 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
99
100 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_scatter_copy_kernel<real>),
101 nblcks, nthrds, 0, strm, (real *) a,
102 (real *) b, (int *) mask, *n, *m);
103
104 HIP_CHECK(hipGetLastError());
105
106 }
107
111 void hip_masked_atomic_reduction(void *a, void *b, void *mask,
112 int *n, int *m, hipStream_t strm) {
113
114 const dim3 nthrds(1024, 1, 1);
115 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
116
117 hipLaunchKernelGGL(HIP_KERNEL_NAME(masked_atomic_reduction_kernel<real>),
118 nblcks, nthrds, 0, strm, (real *) a,
119 (real *) b, (int *) mask, *n, *m);
120
121 HIP_CHECK(hipGetLastError());
122
123 }
124
128 void hip_cfill_mask(void* a, real* c, int* size, void* mask, int* mask_size,
129 hipStream_t strm) {
130
131 const dim3 nthrds(1024, 1, 1);
132 const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
133
134 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_mask_kernel<real>),
135 nblcks, nthrds, 0, strm, (real*)a,
136 *c, *size, (int*)mask, *mask_size);
137
138 HIP_CHECK(hipGetLastError());
139 }
140
144 void hip_rzero(void *a, int *n, hipStream_t strm) {
145 HIP_CHECK(hipMemsetAsync(a, 0, (*n) * sizeof(real), strm));
146 }
147
151 void hip_cmult(void *a, real *c, int *n, hipStream_t strm) {
152
153 const dim3 nthrds(1024, 1, 1);
154 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
155
156 hipLaunchKernelGGL(HIP_KERNEL_NAME(cmult_kernel<real>),
157 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
158 HIP_CHECK(hipGetLastError());
159
160 }
161
165 void hip_cmult2(void *a, void *b, real *c, int *n, hipStream_t strm) {
166
167 const dim3 nthrds(1024, 1, 1);
168 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
169
170 hipLaunchKernelGGL(HIP_KERNEL_NAME(cmult2_kernel<real>),
171 nblcks, nthrds, 0, strm, (real *) a,(real *) b, *c, *n);
172 HIP_CHECK(hipGetLastError());
173
174 }
175
179 void hip_cdiv(void *a, real *c, int *n, hipStream_t strm) {
180
181 const dim3 nthrds(1024, 1, 1);
182 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
183
184 hipLaunchKernelGGL(HIP_KERNEL_NAME(cdiv_kernel<real>),
185 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
186 HIP_CHECK(hipGetLastError());
187
188 }
189
193 void hip_cdiv2(void *a, void *b, real *c, int *n, hipStream_t strm) {
194
195 const dim3 nthrds(1024, 1, 1);
196 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
197
198 hipLaunchKernelGGL(HIP_KERNEL_NAME(cdiv2_kernel<real>),
199 nblcks, nthrds, 0, strm, (real *) a,(real *) b, *c, *n);
200 HIP_CHECK(hipGetLastError());
201
202 }
203
207 void hip_radd(void *a, real *c, int *n, hipStream_t strm) {
208
209 const dim3 nthrds(1024, 1, 1);
210 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
211
212 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd_kernel<real>),
213 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
214 HIP_CHECK(hipGetLastError());
215 }
216
221 void hip_cadd2(void *a, void *b, real *c, int *n, hipStream_t strm) {
222
223 const dim3 nthrds(1024, 1, 1);
224 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
225
226 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd2_kernel<real>),
227 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *c, *n);
228 HIP_CHECK(hipGetLastError());
229 }
230
234 void hip_cfill(void *a, real *c, int *n, hipStream_t strm) {
235
236 const dim3 nthrds(1024, 1, 1);
237 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
238
239 if (*n > 0) {
240 hipLaunchKernelGGL(HIP_KERNEL_NAME(cfill_kernel<real>),
241 nblcks, nthrds, 0, strm, (real *) a, *c, *n);
242 HIP_CHECK(hipGetLastError());
243 }
244
245 }
246
251 void hip_add2(void *a, void *b, 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(add2_kernel<real>),
257 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *n);
258 HIP_CHECK(hipGetLastError());
259
260 }
261
266 void hip_add3(void *a, void *b, void *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(add3_kernel<real>),
272 nblcks, nthrds, 0, strm, (real *) a,
273 (real *) b, (real *) c, *n);
274 HIP_CHECK(hipGetLastError());
275 }
276
281 void hip_add4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm) {
282
283 const dim3 nthrds(1024, 1, 1);
284 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
285
286 hipLaunchKernelGGL(HIP_KERNEL_NAME(add4_kernel<real>),
287 nblcks, nthrds, 0, strm, (real *) a,
288 (real *) b, (real *) c, (real *) d, *n);
289 HIP_CHECK(hipGetLastError());
290 }
291
297 void hip_add2s1(void *a, void *b, real *c1, int *n, hipStream_t strm) {
298
299 const dim3 nthrds(1024, 1, 1);
300 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
301
302 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s1_kernel<real>),
303 nblcks, nthrds, 0, strm, (real *) a,
304 (real *) b, *c1, *n);
305 HIP_CHECK(hipGetLastError());
306 }
307
313 void hip_add2s2(void *a, void *b, real *c1, int *n, hipStream_t strm) {
314
315 const dim3 nthrds(1024, 1, 1);
316 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
317
318 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s2_kernel<real>),
319 nblcks, nthrds, 0, strm, (real *) a, (real *) b,
320 *c1, *n);
321 HIP_CHECK(hipGetLastError());
322 }
323
330 void hip_add2s2_many(void *x, void **p, void *alpha, int *j, int *n,
331 hipStream_t strm) {
332
333 const dim3 nthrds(1024, 1, 1);
334 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
335
336 hipLaunchKernelGGL(HIP_KERNEL_NAME(add2s2_many_kernel<real>),
337 nblcks, nthrds, 0, strm, (real *) x,
338 (const real **) p, (real *) alpha, *j, *n);
339 HIP_CHECK(hipGetLastError());
340
341 }
342
348 void hip_addsqr2s2(void *a, void *b, real *c1, int *n, hipStream_t strm) {
349
350 const dim3 nthrds(1024, 1, 1);
351 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
352
353 hipLaunchKernelGGL(HIP_KERNEL_NAME(addsqr2s2_kernel<real>),
354 nblcks, nthrds, 0, strm, (real *) a,
355 (real *) b, *c1, *n);
356 HIP_CHECK(hipGetLastError());
357 }
358
364 void hip_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n,
365 hipStream_t strm) {
366
367 const dim3 nthrds(1024, 1, 1);
368 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
369
370 hipLaunchKernelGGL(HIP_KERNEL_NAME(add3s2_kernel<real>),
371 nblcks, nthrds, 0, strm, (real *) a,
372 (real *) b, (real *) c, *c1, *c2, *n);
373 HIP_CHECK(hipGetLastError());
374 }
375
381 void hip_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2,
382 real *c3, int *n, hipStream_t strm) {
383
384 const dim3 nthrds(1024, 1, 1);
385 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
386
387 hipLaunchKernelGGL(HIP_KERNEL_NAME(add4s3_kernel<real>),
388 nblcks, nthrds, 0, strm, (real *) a,
389 (real *) b, (real *) c, (real *) d, *c1, *c2, *c3, *n);
390 HIP_CHECK(hipGetLastError());
391 }
392
398 void hip_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1,
399 real *c2, real *c3, real *c4, int *n, hipStream_t strm) {
400
401 const dim3 nthrds(1024, 1, 1);
402 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
403
404 hipLaunchKernelGGL(HIP_KERNEL_NAME(add5s4_kernel<real>),
405 nblcks, nthrds, 0, strm, (real *) a,
406 (real *) b, (real *) c, (real *) d, (real *) e,
407 *c1, *c2, *c3, *c4, *n);
408 HIP_CHECK(hipGetLastError());
409 }
410
415 void hip_invcol1(void *a, int *n, hipStream_t strm) {
416
417 const dim3 nthrds(1024, 1, 1);
418 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
419
420 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol1_kernel<real>),
421 nblcks, nthrds, 0, strm, (real *) a, *n);
422 HIP_CHECK(hipGetLastError());
423 }
424
429 void hip_invcol2(void *a, void *b, int *n, hipStream_t strm) {
430
431 const dim3 nthrds(1024, 1, 1);
432 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
433
434 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol2_kernel<real>),
435 nblcks, nthrds, 0, strm, (real *) a, (real *) b, *n);
436 HIP_CHECK(hipGetLastError());
437 }
438
443 void hip_invcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
444
445 const dim3 nthrds(1024, 1, 1);
446 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
447
448 hipLaunchKernelGGL(HIP_KERNEL_NAME(invcol3_kernel<real>),
449 nblcks, nthrds, 0, strm, (real *) a,
450 (real *) b, (real *) c, *n);
451 HIP_CHECK(hipGetLastError());
452 }
453
458 void hip_col2(void *a, void *b, int *n, hipStream_t strm) {
459
460 const dim3 nthrds(1024, 1, 1);
461 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
462
463 hipLaunchKernelGGL(HIP_KERNEL_NAME(col2_kernel<real>),
464 nblcks, nthrds, 0, strm, (real *) a,
465 (real *) b, *n);
466 HIP_CHECK(hipGetLastError());
467 }
468
473 void hip_col3(void *a, void *b, void *c, int *n, hipStream_t strm) {
474
475 const dim3 nthrds(1024, 1, 1);
476 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
477
478 hipLaunchKernelGGL(HIP_KERNEL_NAME(col3_kernel<real>),
479 nblcks, nthrds, 0, strm, (real *) a,
480 (real *) b, (real *) c, *n);
481 HIP_CHECK(hipGetLastError());
482 }
483
488 void hip_subcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
489
490 const dim3 nthrds(1024, 1, 1);
491 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
492
493 hipLaunchKernelGGL(HIP_KERNEL_NAME(subcol3_kernel<real>),
494 nblcks, nthrds, 0, strm, (real *) a,
495 (real *) b, (real *) c, *n);
496 HIP_CHECK(hipGetLastError());
497 }
498
503 void hip_sub2(void *a, void *b, int *n, hipStream_t strm) {
504
505 const dim3 nthrds(1024, 1, 1);
506 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
507
508 hipLaunchKernelGGL(HIP_KERNEL_NAME(sub2_kernel<real>),
509 nblcks, nthrds, 0, strm, (real *) a,
510 (real *) b, *n);
511 HIP_CHECK(hipGetLastError());
512 }
513
518 void hip_sub3(void *a, void *b, void *c, int *n, hipStream_t strm) {
519
520 const dim3 nthrds(1024, 1, 1);
521 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
522
523 hipLaunchKernelGGL(HIP_KERNEL_NAME(sub3_kernel<real>),
524 nblcks, nthrds, 0, strm, (real *) a,
525 (real *) b, (real *) c, *n);
526 HIP_CHECK(hipGetLastError());
527 }
528
533 void hip_addcol3(void *a, void *b, void *c, int *n, hipStream_t strm) {
534
535 const dim3 nthrds(1024, 1, 1);
536 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
537
538 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol3_kernel<real>),
539 nblcks, nthrds, 0, strm, (real *) a,
540 (real *) b, (real *) c, *n);
541 HIP_CHECK(hipGetLastError());
542 }
543
548 void hip_addcol4(void *a, void *b, void *c, void *d, int *n,
549 hipStream_t strm) {
550
551 const dim3 nthrds(1024, 1, 1);
552 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
553
554 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol4_kernel<real>),
555 nblcks, nthrds, 0, strm, (real *) a,
556 (real *) b, (real *) c, (real *) d, *n);
557 HIP_CHECK(hipGetLastError());
558 }
559
564 void hip_addcol3s2(void *a, void *b, void *c, real *s, int *n,
565 hipStream_t strm) {
566
567 const dim3 nthrds(1024, 1, 1);
568 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
569
570 hipLaunchKernelGGL(HIP_KERNEL_NAME(addcol3s2_kernel<real>),
571 nblcks, nthrds, 0, strm, (real *) a,
572 (real *) b, (real *) c, *s, *n);
573 HIP_CHECK(hipGetLastError());
574 }
575
580 void hip_vdot3(void *dot, void *u1, void *u2, void *u3,
581 void *v1, void *v2, void *v3, int *n,
582 hipStream_t strm) {
583
584 const dim3 nthrds(1024, 1, 1);
585 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
586
587 hipLaunchKernelGGL(HIP_KERNEL_NAME(vdot3_kernel<real>),
588 nblcks, nthrds, 0, strm,
589 (real *) dot, (real *) u1, (real *) u2, (real *) u3,
590 (real *) v1, (real *) v2, (real *) v3, *n);
591 HIP_CHECK(hipGetLastError());
592 }
593
598 void hip_vcross(void *u1, void *u2, void *u3,
599 void *v1, void *v2, void *v3,
600 void *w1, void *w2, void *w3,
601 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(vcross_kernel<real>),
607 nblcks, nthrds, 0, strm,
608 (real *) u1, (real *) u2, (real *) u3,
609 (real *) v1, (real *) v2, (real *) v3,
610 (real *) w1, (real *) w2, (real *) w3, *n);
611 HIP_CHECK(hipGetLastError());
612 }
613
614
615 /*
616 * Reduction buffer
617 */
618 int red_s = 0;
619 real * bufred = NULL;
620 real * bufred_d = NULL;
621
622 void hip_redbuf_check_alloc(int nb) {
623 if ( nb >= red_s) {
624 red_s = nb+1;
625 if (bufred != NULL) {
626 HIP_CHECK(hipHostFree(bufred));
627 HIP_CHECK(hipFree(bufred_d));
628 }
629 HIP_CHECK(hipHostMalloc(&bufred,red_s*sizeof(real),hipHostMallocDefault));
630 HIP_CHECK(hipMalloc(&bufred_d, red_s*sizeof(real)));
631 }
632 }
633
637 void hip_global_reduce_add(real * bufred, void * bufred_d, int n,
638 const hipStream_t stream) {
639 #ifdef HAVE_RCCL
640 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
641 DEVICE_NCCL_SUM, stream);
642 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
643 hipMemcpyDeviceToHost, stream));
644 HIP_CHECK(hipStreamSynchronize(stream));
645 #elif HAVE_DEVICE_MPI
646 HIP_CHECK(hipStreamSynchronize(stream));
647 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_SUM);
648 #else
649 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, n*sizeof(real),
650 hipMemcpyDeviceToHost, stream));
651 HIP_CHECK(hipStreamSynchronize(stream));
652 #endif
653 }
654
659 real hip_vlsc3(void *u, void *v, void *w, int *n, hipStream_t stream) {
660
661 const dim3 nthrds(1024, 1, 1);
662 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
663 const int nb = ((*n) + 1024 - 1)/ 1024;
664
665 hip_redbuf_check_alloc(nb);
666
667
668 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_kernel<real>),
669 nblcks, nthrds, 0, stream,
670 (real *) u, (real *) v,
671 (real *) w, bufred_d, *n);
672 HIP_CHECK(hipGetLastError());
673 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
674 1, 1024, 0, stream, bufred_d, nb);
675 HIP_CHECK(hipGetLastError());
676
677 HIP_CHECK(hipMemcpyAsync(bufred, bufred_d, sizeof(real),
678 hipMemcpyDeviceToHost, stream));
679 HIP_CHECK(hipStreamSynchronize(stream));
680
681 return bufred[0];
682 }
683
688 real hip_glsc3(void *a, void *b, void *c, int *n, hipStream_t stream) {
689
690 const dim3 nthrds(1024, 1, 1);
691 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
692 const int nb = ((*n) + 1024 - 1)/ 1024;
693
694 hip_redbuf_check_alloc(nb);
695
696 if (*n > 0) {
697 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_kernel<real>),
698 nblcks, nthrds, 0, stream,
699 (real *) a, (real *) b,
700 (real *) c, bufred_d, *n);
701 HIP_CHECK(hipGetLastError());
702 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
703 1, 1024, 0, stream, bufred_d, nb);
704 HIP_CHECK(hipGetLastError());
705 }
706 else {
707 hip_rzero(bufred_d, &red_s, stream);
708 }
709 hip_global_reduce_add(bufred, bufred_d, 1, stream);
710
711 return bufred[0];
712 }
713
718 void hip_glsc3_many(real *h, void * w, void *v,void *mult, int *j, int *n,
719 hipStream_t stream){
720 int pow2 = 1;
721 while(pow2 < (*j)){
722 pow2 = 2*pow2;
723 }
724 const int nt = 1024/pow2;
725 const dim3 nthrds(pow2, nt, 1);
726 const dim3 nblcks(((*n)+nt - 1)/nt, 1, 1);
727 const dim3 nthrds_red(1024,1,1);
728 const dim3 nblcks_red( (*j),1,1);
729 const int nb = ((*n) + nt - 1)/nt;
730
731 hip_redbuf_check_alloc((*j)*nb);
732
733 if (*n > 0) {
734 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_many_kernel<real>),
735 nblcks, nthrds, 0, stream,
736 (const real *) w, (const real **) v,
737 (const real *)mult, bufred_d, *j, *n);
738 HIP_CHECK(hipGetLastError());
739
740 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc3_reduce_kernel<real>),
741 nblcks_red, nthrds_red, 0, stream,
742 bufred_d, nb, *j);
743 HIP_CHECK(hipGetLastError());
744 }
745 else {
746 hip_rzero(bufred_d, &red_s, stream);
747 }
748 hip_global_reduce_add(h, bufred_d, (*j), stream);
749 }
750
755 real hip_glsc2(void *a, void *b, int *n, hipStream_t stream) {
756
757 const dim3 nthrds(1024, 1, 1);
758 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
759 const int nb = ((*n) + 1024 - 1)/ 1024;
760
761
762 hip_redbuf_check_alloc(nb);
763
764 if( *n > 0) {
765 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsc2_kernel<real>),
766 nblcks, nthrds, 0, stream,
767 (real *) a, (real *) b, bufred_d, *n);
768 HIP_CHECK(hipGetLastError());
769 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
770 1, 1024, 0, stream, bufred_d, nb);
771 HIP_CHECK(hipGetLastError());
772 }
773 else {
774 hip_rzero(bufred_d, &red_s, stream);
775 }
776 hip_global_reduce_add(bufred, bufred_d, 1, stream);
777
778 return bufred[0];
779 }
780
785 real hip_glsubnorm2(void* a, void* b, int* n, hipStream_t stream) {
786
787 const dim3 nthrds(1024, 1, 1);
788 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
789 const int nb = ((*n) + 1024 - 1) / 1024;
790
791 hip_redbuf_check_alloc(nb);
792
793 if (*n > 0) {
794 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsubnorm2_kernel<real>),
795 nblcks, nthrds, 0, stream,
796 (real*)a, (real*)b, bufred_d, *n);
797 HIP_CHECK(hipGetLastError());
798 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
799 1, 1024, 0, stream, bufred_d, nb);
800 HIP_CHECK(hipGetLastError());
801 }
802 else {
803 hip_rzero(bufred_d, &red_s, stream);
804 }
805 hip_global_reduce_add(bufred, bufred_d, 1, stream);
806
807 return bufred[0];
808 }
809
814 real hip_glsum(void *a, int *n, hipStream_t stream) {
815 const dim3 nthrds(1024, 1, 1);
816 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
817 const int nb = ((*n) + 1024 - 1)/ 1024;
818
819 hip_redbuf_check_alloc(nb);
820 if( *n > 0) {
821 hipLaunchKernelGGL(HIP_KERNEL_NAME(glsum_kernel<real>),
822 nblcks, nthrds, 0, stream,
823 (real *) a, bufred_d, *n);
824 HIP_CHECK(hipGetLastError());
825 hipLaunchKernelGGL(HIP_KERNEL_NAME(reduce_kernel<real>),
826 1, 1024, 0, stream, bufred_d, nb);
827 HIP_CHECK(hipGetLastError());
828 }
829 else {
830 hip_rzero(bufred_d, &red_s, stream);
831 }
832
833 hip_global_reduce_add(bufred, bufred_d, 1, stream);
834
835 return bufred[0];
836 }
837
842 void hip_absval(void *a, int *n, hipStream_t stream) {
843
844 const dim3 nthrds(1024, 1, 1);
845 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
846
847 hipLaunchKernelGGL(HIP_KERNEL_NAME(absval_kernel<real>),
848 nblcks, nthrds, 0, stream, (real *) a, *n);
849 HIP_CHECK(hipGetLastError());
850
851}
852
853 // ======================================================================== //
854 // Point-wise operations.
855
860 void hip_pwmax_vec2(void* a, void* b, int* n, hipStream_t stream) {
861
862 const dim3 nthrds(1024, 1, 1);
863 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
864
865 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_vec2_kernel<real>),
866 nblcks, nthrds, 0, stream, (real*)a, (real*)b, *n);
867 HIP_CHECK(hipGetLastError());
868}
869
874 void hip_pwmax_vec3(void *a, void *b, void *c, int *n, hipStream_t stream) {
875
876 const dim3 nthrds(1024, 1, 1);
877 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
878
879 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_vec3_kernel<real>),
880 nblcks, nthrds, 0, stream,
881 (real *)a, (real *)b, (real *)c, *n);
882 HIP_CHECK(hipGetLastError());
883 }
884
889 void hip_pwmax_sca2(void *a, real *c, int *n, hipStream_t stream) {
890
891 const dim3 nthrds(1024, 1, 1);
892 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
893
894 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_sca2_kernel<real>),
895 nblcks, nthrds, 0, stream, (real *)a, *c, *n);
896 HIP_CHECK(hipGetLastError());
897 }
898
903 void hip_pwmax_sca3(void *a, void *b, real *c, int *n, hipStream_t stream) {
904
905 const dim3 nthrds(1024, 1, 1);
906 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
907
908 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmax_sca3_kernel<real>),
909 nblcks, nthrds, 0, stream,
910 (real *)a, (real *)b, *c, *n);
911 HIP_CHECK(hipGetLastError());
912 }
913
918 void hip_pwmin_vec2(void *a, void *b, int *n, hipStream_t stream) {
919
920 const dim3 nthrds(1024, 1, 1);
921 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
922
923 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_vec2_kernel<real>),
924 nblcks, nthrds, 0, stream,
925 (real *)a, (real *)b, *n);
926 HIP_CHECK(hipGetLastError());
927 }
928
933 void hip_pwmin_vec3(void *a, void *b, void *c, int *n, hipStream_t stream) {
934
935 const dim3 nthrds(1024, 1, 1);
936 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
937
938 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_vec3_kernel<real>),
939 nblcks, nthrds, 0, stream,
940 (real *)a, (real *)b, (real *)c, *n);
941 HIP_CHECK(hipGetLastError());
942 }
943
948 void hip_pwmin_sca2(void *a, real *c, int *n, hipStream_t stream) {
949
950 const dim3 nthrds(1024, 1, 1);
951 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
952
953 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_sca2_kernel<real>),
954 nblcks, nthrds, 0, stream,
955 (real *)a, *c, *n);
956 HIP_CHECK(hipGetLastError());
957 }
958
963 void hip_pwmin_sca3(void *a, void *b, real *c, int *n, hipStream_t stream) {
964
965 const dim3 nthrds(1024, 1, 1);
966 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
967
968 hipLaunchKernelGGL(HIP_KERNEL_NAME(pwmin_sca3_kernel<real>),
969 nblcks, nthrds, 0, stream,
970 (real *)a, (real *)b, *c, *n);
971 HIP_CHECK(hipGetLastError());
972 }
973
974 // ======================================================================== //
975
979 void hip_iadd(void *a, int *c, int *n, hipStream_t stream) {
980
981 const dim3 nthrds(1024, 1, 1);
982 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
983
984 hipLaunchKernelGGL(HIP_KERNEL_NAME(cadd_kernel<int>),
985 nblcks, nthrds, 0, stream,
986 (int *) a, *c, *n);
987 HIP_CHECK(hipGetLastError());
988 }
989
990} /* extern "C" */
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
dirichlet_apply_scalar_kernel
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
Definition dirichlet_kernel.h:42
device_config.h
real
double real
Definition device_config.h:12
device_mpi_op.h
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_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_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:580
hip_global_reduce_add
void hip_global_reduce_add(real *bufred, void *bufred_d, int n, const hipStream_t stream)
Definition math.hip:637
hip_col3
void hip_col3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:473
hip_addsqr2s2
void hip_addsqr2s2(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:348
hip_rzero
void hip_rzero(void *a, int *n, hipStream_t strm)
Definition math.hip:144
hip_cfill_mask
void hip_cfill_mask(void *a, real *c, int *size, void *mask, int *mask_size, hipStream_t strm)
Definition math.hip:128
hip_pwmax_sca3
void hip_pwmax_sca3(void *a, void *b, real *c, int *n, hipStream_t stream)
Definition math.hip:903
hip_invcol3
void hip_invcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:443
hip_absval
void hip_absval(void *a, int *n, hipStream_t stream)
Definition math.hip:842
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:598
hip_glsubnorm2
real hip_glsubnorm2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:785
hip_addcol3
void hip_addcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:533
hip_add3s2
void hip_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n, hipStream_t strm)
Definition math.hip:364
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:398
hip_add2s2
void hip_add2s2(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:313
hip_glsum
real hip_glsum(void *a, int *n, hipStream_t stream)
Definition math.hip:814
hip_add4
void hip_add4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm)
Definition math.hip:281
hip_cadd2
void hip_cadd2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:221
hip_vlsc3
real hip_vlsc3(void *u, void *v, void *w, int *n, hipStream_t stream)
Definition math.hip:659
hip_cdiv2
void hip_cdiv2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:193
hip_pwmin_sca2
void hip_pwmin_sca2(void *a, real *c, int *n, hipStream_t stream)
Definition math.hip:948
hip_cmult2
void hip_cmult2(void *a, void *b, real *c, int *n, hipStream_t strm)
Definition math.hip:165
hip_glsc3
real hip_glsc3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:688
hip_invcol1
void hip_invcol1(void *a, int *n, hipStream_t strm)
Definition math.hip:415
hip_invcol2
void hip_invcol2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:429
hip_masked_copy
void hip_masked_copy(void *a, void *b, void *mask, int *n, int *m, hipStream_t strm)
Definition math.hip:61
hip_cdiv
void hip_cdiv(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:179
hip_sub2
void hip_sub2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:503
hip_glsc2
real hip_glsc2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:755
bufred
real * bufred
Definition math.hip:619
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:381
hip_col2
void hip_col2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:458
hip_pwmin_vec3
void hip_pwmin_vec3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:933
hip_addcol3s2
void hip_addcol3s2(void *a, void *b, void *c, real *s, int *n, hipStream_t strm)
Definition math.hip:564
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:78
hip_glsc3_many
void hip_glsc3_many(real *h, void *w, void *v, void *mult, int *j, int *n, hipStream_t stream)
Definition math.hip:718
red_s
int red_s
Definition math.hip:618
hip_cfill
void hip_cfill(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:234
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:111
hip_add3
void hip_add3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:266
hip_pwmax_vec3
void hip_pwmax_vec3(void *a, void *b, void *c, int *n, hipStream_t stream)
Definition math.hip:874
hip_add2
void hip_add2(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:251
hip_copy
void hip_copy(void *a, void *b, int *n, hipStream_t strm)
Definition math.hip:53
hip_subcol3
void hip_subcol3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:488
hip_pwmax_sca2
void hip_pwmax_sca2(void *a, real *c, int *n, hipStream_t stream)
Definition math.hip:889
hip_iadd
void hip_iadd(void *a, int *c, int *n, hipStream_t stream)
Definition math.hip:979
hip_pwmin_sca3
void hip_pwmin_sca3(void *a, void *b, real *c, int *n, hipStream_t stream)
Definition math.hip:963
hip_add2s1
void hip_add2s1(void *a, void *b, real *c1, int *n, hipStream_t strm)
Definition math.hip:297
hip_radd
void hip_radd(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:207
hip_redbuf_check_alloc
void hip_redbuf_check_alloc(int nb)
Definition math.hip:622
hip_add2s2_many
void hip_add2s2_many(void *x, void **p, void *alpha, int *j, int *n, hipStream_t strm)
Definition math.hip:330
hip_addcol4
void hip_addcol4(void *a, void *b, void *c, void *d, int *n, hipStream_t strm)
Definition math.hip:548
hip_pwmax_vec2
void hip_pwmax_vec2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:860
hip_pwmin_vec2
void hip_pwmin_vec2(void *a, void *b, int *n, hipStream_t stream)
Definition math.hip:918
hip_sub3
void hip_sub3(void *a, void *b, void *c, int *n, hipStream_t strm)
Definition math.hip:518
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:94
hip_cmult
void hip_cmult(void *a, real *c, int *n, hipStream_t strm)
Definition math.hip:151
bufred_d
real * bufred_d
Definition math.hip:620
mask
Object for handling masks in Neko.
Definition mask.f90:34