Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
math.cu
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 "math_kernel.h"
36#include <device/device_config.h>
37#include <device/cuda/check.h>
38#include <device/cuda/buffer.h>
39#include <stdio.h>
40#include <stdlib.h>
41
42#ifdef HAVE_NVSHMEM
43#include <nvshmem.h>
44#include <nvshmemx.h>
45#endif
46
47extern "C" {
48
49#include <math/bcknd/device/device_mpi_reduce.h>
50#include <math/bcknd/device/device_mpi_op.h>
51
52#ifdef HAVE_NCCL
53#include <math/bcknd/device/device_nccl_reduce.h>
54#include <math/bcknd/device/device_nccl_op.h>
55#endif
56
60 void cuda_copy(void *a, void *b, int *n, cudaStream_t strm) {
61 CUDA_CHECK(cudaMemcpyAsync(a, b, (*n) * sizeof(real),
62 cudaMemcpyDeviceToDevice, strm));
63 }
64
69 void cuda_masked_copy_0(void *a, void *b, void *mask,
70 int *n, int *m, cudaStream_t strm) {
71
72 const dim3 nthrds(1024, 1, 1);
73 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
74
75 masked_copy_kernel_0<real><<<nblcks, nthrds, 0, strm>>>
76 ((real *) a, (real*) b,(int*) mask, *n, *m);
77 CUDA_CHECK(cudaGetLastError());
78
79 }
80
85 void cuda_masked_copy_aligned(void *a, void *b, void *mask,
86 int *n, int *m, cudaStream_t strm) {
87
88 const dim3 nthrds(1024, 1, 1);
89 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
90
91 masked_copy_kernel_aligned<real><<<nblcks, nthrds, 0, strm>>>
92 ((real *) a, (real*) b,(int*) mask, *n, *m);
93 CUDA_CHECK(cudaGetLastError());
94
95 }
96
100 void cuda_masked_gather_copy(void *a, void *b, void *mask,
101 int *n, int *m, cudaStream_t strm) {
102
103 const dim3 nthrds(1024, 1, 1);
104 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
105
106 masked_gather_copy_kernel<real><<<nblcks, nthrds, 0, strm>>>
107 ((real *) a, (real*) b,(int*) mask, *n, *m);
108 CUDA_CHECK(cudaGetLastError());
109
110 }
111
115 void cuda_masked_gather_copy_aligned(void *a, void *b, void *mask,
116 int *n, int *m, cudaStream_t strm) {
117
118 const dim3 nthrds(1024, 1, 1);
119 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
120
121 masked_gather_copy_aligned_kernel<real><<<nblcks, nthrds, 0, strm>>>
122 ((real *) a, (real*) b,(int*) mask, *n, *m);
123 CUDA_CHECK(cudaGetLastError());
124
125 }
126
130 void cuda_face_masked_gather_copy(void *a, void *b, void *mask,
131 void *facet, int *n1, int *n2, int *lx,
132 int *ly, int *lz, int *m,
133 cudaStream_t strm) {
134
135 const dim3 nthrds(1024, 1, 1);
136 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
137
138 face_masked_gather_copy_kernel<real><<<nblcks, nthrds, 0, strm>>>
139 ((real *) a, (real *) b, (int *) mask, (int *) facet, *n1, *n2, *lx,
140 *ly, *lz, *m);
141 CUDA_CHECK(cudaGetLastError());
142
143 }
144
148 void cuda_masked_atomic_reduction(void *a, void *b, void *mask,
149 int *n, int *m, cudaStream_t strm) {
150
151 const dim3 nthrds(MASKED_ATOMIC_RED_BLOCK, 1, 1);
152 const dim3 nblcks(((*m) + MASKED_ATOMIC_RED_BLOCK - 1) /
153 MASKED_ATOMIC_RED_BLOCK, 1, 1);
154
155 masked_atomic_reduction_kernel<real><<<nblcks, nthrds, 0, strm>>>
156 ((real *) a, (real *) b, (int *) mask, *n, *m);
157 CUDA_CHECK(cudaGetLastError());
158
159 }
163 void cuda_masked_scatter_copy(void *a, void *b, void *mask,
164 int *n, int *m, cudaStream_t strm) {
165
166 const dim3 nthrds(1024, 1, 1);
167 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
168
169 masked_scatter_copy_kernel<real><<<nblcks, nthrds, 0, strm>>>
170 ((real *) a, (real*) b,(int*) mask, *n, *m);
171 CUDA_CHECK(cudaGetLastError());
172 }
173
177 void cuda_masked_scatter_copy_aligned(void *a, void *b, void *mask,
178 int *n, int *m, cudaStream_t strm) {
179
180 const dim3 nthrds(1024, 1, 1);
181 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
182
183 masked_scatter_copy_aligned_kernel<real><<<nblcks, nthrds, 0, strm>>>
184 ((real *) a, (real*) b,(int*) mask, *n, *m);
185 CUDA_CHECK(cudaGetLastError());
186 }
187
188
192 void cuda_cfill_mask(void* a, real* c, int* size, int* mask, int* mask_size,
193 cudaStream_t strm) {
194
195 const dim3 nthrds(1024, 1, 1);
196 const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
197
198 cfill_mask_kernel<real><<<nblcks, nthrds, 0, strm>>>
199 ((real*)a, *c, *size, mask, *mask_size);
200 CUDA_CHECK(cudaGetLastError());
201 }
202
206 void cuda_rzero(void *a, int *n, cudaStream_t strm) {
207 CUDA_CHECK(cudaMemsetAsync(a, 0, (*n) * sizeof(real), strm));
208 }
209
213 void cuda_cmult(void *a, real *c, int *n, cudaStream_t strm) {
214
215 const dim3 nthrds(1024, 1, 1);
216 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
217
218 cmult_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
219 CUDA_CHECK(cudaGetLastError());
220
221 }
222
226 void cuda_cmult2(void *a, void *b, real *c, int *n, cudaStream_t strm) {
227
228 const dim3 nthrds(1024, 1, 1);
229 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
230
231 cmult2_kernel<real><<<nblcks, nthrds, 0, strm>>>
232 ((real *) a, (real *) b, *c, *n);
233 CUDA_CHECK(cudaGetLastError());
234
235 }
236
240 void cuda_cdiv(void *a, real *c, int *n, cudaStream_t strm) {
241
242 const dim3 nthrds(1024, 1, 1);
243 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
244
245 cdiv_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
246 CUDA_CHECK(cudaGetLastError());
247
248 }
249
253 void cuda_cdiv2(void *a, void *b, real *c, int *n, cudaStream_t strm) {
254
255 const dim3 nthrds(1024, 1, 1);
256 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
257
258 cdiv2_kernel<real><<<nblcks, nthrds, 0, strm>>>
259 ((real *) a, (real *) b, *c, *n);
260 CUDA_CHECK(cudaGetLastError());
261
262 }
263
267 void cuda_radd(void *a, real *c, int *n, cudaStream_t strm) {
268
269 const dim3 nthrds(1024, 1, 1);
270 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
271
272 cadd_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
273 CUDA_CHECK(cudaGetLastError());
274
275 }
276
281 void cuda_cadd2(void *a, void *b, real *c, int *n, cudaStream_t strm) {
282
283 const dim3 nthrds(1024, 1, 1);
284 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
285
286 cadd2_kernel<real><<<nblcks, nthrds, 0, strm>>>
287 ((real *) a, (real *) b, *c, *n);
288 CUDA_CHECK(cudaGetLastError());
289
290 }
291
296 void cuda_cwrap(void *a, real *min_val, real *max_val, int *n,
297 cudaStream_t strm) {
298
299 const dim3 nthrds(1024, 1, 1);
300 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
301
302 cwrap_kernel<real><<<nblcks, nthrds, 0, strm>>>
303 ((real *) a, *min_val, *max_val, *n);
304 CUDA_CHECK(cudaGetLastError());
305
306 }
307
311 void cuda_sqrt_inplace(void *a, int *n, cudaStream_t strm) {
312
313 const dim3 nthrds(1024, 1, 1);
314 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
315
316 sqrt_inplace_kernel<real><<<nblcks, nthrds, 0, strm>>>
317 ((real *) a, *n);
318 CUDA_CHECK(cudaGetLastError());
319
320 }
321
325 void cuda_power(void *ap, void *a, real *p, int *n, cudaStream_t strm) {
326
327 const dim3 nthrds(1024, 1, 1);
328 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
329
330 power_kernel<real><<<nblcks, nthrds, 0, strm>>>
331 ((real *) ap, (real *) a, *p, *n);
332 CUDA_CHECK(cudaGetLastError());
333
334 }
335
339 void cuda_cfill(void *a, real *c, int *n, cudaStream_t strm) {
340
341 const dim3 nthrds(1024, 1, 1);
342 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
343
344 if (*n > 0){
345 cfill_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
346 CUDA_CHECK(cudaGetLastError());
347 }
348
349 }
350
355 void cuda_add2(void *a, void *b, int *n, cudaStream_t strm) {
356
357 const dim3 nthrds(1024, 1, 1);
358 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
359
360 add2_kernel<real><<<nblcks, nthrds, 0, strm>>>
361 ((real *) a, (real *) b, *n);
362 CUDA_CHECK(cudaGetLastError());
363
364 }
365
370 void cuda_add3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
371
372 const dim3 nthrds(1024, 1, 1);
373 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
374
375 add3_kernel<real><<<nblcks, nthrds, 0, strm>>>
376 ((real *) a, (real *) b, (real *) c, *n);
377 CUDA_CHECK(cudaGetLastError());
378 }
379
384 void cuda_add4(void *a, void *b, void *c, void *d, int *n,
385 cudaStream_t strm) {
386
387 const dim3 nthrds(1024, 1, 1);
388 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
389
390 add4_kernel<real><<<nblcks, nthrds, 0, strm>>>
391 ((real *) a, (real *) b, (real *) c, (real *) d, *n);
392 CUDA_CHECK(cudaGetLastError());
393
394 }
400 void cuda_add2s1(void *a, void *b, real *c1, int *n, cudaStream_t strm) {
401
402 const dim3 nthrds(1024, 1, 1);
403 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
404
405 add2s1_kernel<real><<<nblcks, nthrds, 0, strm>>>
406 ((real *) a, (real *) b, *c1, *n);
407 CUDA_CHECK(cudaGetLastError());
408
409 }
410
416 void cuda_add2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm) {
417
418 const dim3 nthrds(1024, 1, 1);
419 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
420
421 add2s2_kernel<real><<<nblcks, nthrds, 0, strm>>>
422 ((real *) a, (real *) b, *c1, *n);
423 CUDA_CHECK(cudaGetLastError());
424
425 }
426
433 void cuda_add2s2_many(void *x, void **p, void *alpha, int *j, int *n,
434 cudaStream_t strm) {
435
436 const dim3 nthrds(1024, 1, 1);
437 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
438
439 add2s2_many_kernel<real><<<nblcks, nthrds, 0, strm>>>
440 ((real *) x, (const real **) p, (real *) alpha, *j, *n);
441 CUDA_CHECK(cudaGetLastError());
442
443 }
444
450 void cuda_addsqr2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm) {
451
452 const dim3 nthrds(1024, 1, 1);
453 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
454
455 addsqr2s2_kernel<real><<<nblcks, nthrds, 0,strm>>>
456 ((real *) a, (real *) b, *c1, *n);
457 CUDA_CHECK(cudaGetLastError());
458
459 }
460
466 void cuda_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n,
467 cudaStream_t strm) {
468
469 const dim3 nthrds(1024, 1, 1);
470 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
471
472 add3s2_kernel<real><<<nblcks, nthrds, 0, strm>>>
473 ((real *) a, (real *) b, (real *) c, *c1, *c2, *n);
474 CUDA_CHECK(cudaGetLastError());
475
476 }
477
483 void cuda_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2,
484 real *c3, int *n, cudaStream_t strm) {
485
486 const dim3 nthrds(1024, 1, 1);
487 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
488
489 add4s3_kernel<real><<<nblcks, nthrds, 0, strm>>>
490 ((real *) a, (real *) b, (real *) c, (real *) d, *c1, *c2, *c3, *n);
491 CUDA_CHECK(cudaGetLastError());
492
493 }
494
500 void cuda_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1,
501 real *c2, real *c3, real *c4, int *n, cudaStream_t strm) {
502
503 const dim3 nthrds(1024, 1, 1);
504 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
505
506 add5s4_kernel<real><<<nblcks, nthrds, 0, strm>>>
507 ((real *) a, (real *) b, (real *) c, (real *) d, (real *) e,
508 *c1, *c2, *c3, *c4, *n);
509 CUDA_CHECK(cudaGetLastError());
510
511 }
512
517 void cuda_invcol1(void *a, int *n, cudaStream_t strm) {
518
519 const dim3 nthrds(1024, 1, 1);
520 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
521
522 invcol1_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *n);
523 CUDA_CHECK(cudaGetLastError());
524 }
525
530 void cuda_invcol2(void *a, void *b, int *n, cudaStream_t strm) {
531
532 const dim3 nthrds(1024, 1, 1);
533 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
534
535 invcol2_kernel<real><<<nblcks, nthrds, 0, strm>>>
536 ((real *) a, (real *) b, *n);
537 CUDA_CHECK(cudaGetLastError());
538 }
539
544 void cuda_invcol3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
545
546 const dim3 nthrds(1024, 1, 1);
547 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
548
549 invcol3_kernel<real><<<nblcks, nthrds, 0, strm>>>
550 ((real *) a, (real *) b, (real *) c, *n);
551 CUDA_CHECK(cudaGetLastError());
552 }
553
558 void cuda_col2(void *a, void *b, int *n, cudaStream_t strm) {
559
560 const dim3 nthrds(1024, 1, 1);
561 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
562
563 col2_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, (real *) b, *n);
564 CUDA_CHECK(cudaGetLastError());
565 }
566
571 void cuda_col3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
572
573 const dim3 nthrds(1024, 1, 1);
574 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
575
576 col3_kernel<real><<<nblcks, nthrds, 0, strm>>>
577 ((real *) a, (real *) b, (real *) c, *n);
578 CUDA_CHECK(cudaGetLastError());
579 }
580
585 void cuda_subcol3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
586
587 const dim3 nthrds(1024, 1, 1);
588 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
589
590 subcol3_kernel<real><<<nblcks, nthrds, 0, strm>>>
591 ((real *) a, (real *) b, (real *) c, *n);
592 CUDA_CHECK(cudaGetLastError());
593 }
594
595
600 void cuda_sub2(void *a, void *b, int *n, cudaStream_t strm) {
601
602 const dim3 nthrds(1024, 1, 1);
603 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
604
605 sub2_kernel<real><<<nblcks, nthrds, 0, strm>>>
606 ((real *) a, (real *) b, *n);
607 CUDA_CHECK(cudaGetLastError());
608 }
609
614 void cuda_sub3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
615
616 const dim3 nthrds(1024, 1, 1);
617 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
618
619 sub3_kernel<real><<<nblcks, nthrds, 0, strm>>>
620 ((real *) a, (real *) b, (real *) c, *n);
621 CUDA_CHECK(cudaGetLastError());
622 }
623
628 void cuda_addcol3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
629
630 const dim3 nthrds(1024, 1, 1);
631 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
632
633 addcol3_kernel<real><<<nblcks, nthrds, 0, strm>>>
634 ((real *) a, (real *) b, (real *) c, *n);
635 CUDA_CHECK(cudaGetLastError());
636 }
637
642 void cuda_addcol4(void *a, void *b, void *c, void *d, int *n,
643 cudaStream_t strm) {
644
645 const dim3 nthrds(1024, 1, 1);
646 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
647
648 addcol4_kernel<real><<<nblcks, nthrds, 0, strm>>>
649 ((real *) a, (real *) b, (real *) c, (real *) d, *n);
650 CUDA_CHECK(cudaGetLastError());
651 }
652
657 void cuda_addcol3s2(void *a, void *b, void *c, real *s, int *n,
658 cudaStream_t strm) {
659
660 const dim3 nthrds(1024, 1, 1);
661 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
662
663 addcol3s2_kernel<real><<<nblcks, nthrds, 0, strm>>>
664 ((real *) a, (real *) b, (real *) c, *s, *n);
665 CUDA_CHECK(cudaGetLastError());
666 }
667
672 void cuda_vdot3(void *dot, void *u1, void *u2, void *u3,
673 void *v1, void *v2, void *v3, int *n,
674 cudaStream_t strm) {
675
676 const dim3 nthrds(1024, 1, 1);
677 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
678
679 vdot3_kernel<real><<<nblcks, nthrds, 0, strm>>>
680 ((real *) dot, (real *) u1, (real *) u2, (real *) u3,
681 (real *) v1, (real *) v2, (real *) v3, *n);
682 CUDA_CHECK(cudaGetLastError());
683 }
684
689 void cuda_vcross(void *u1, void *u2, void *u3,
690 void *v1, void *v2, void *v3,
691 void *w1, void *w2, void *w3,
692 int *n, cudaStream_t strm) {
693
694 const dim3 nthrds(1024, 1, 1);
695 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
696
697 vcross_kernel<real><<<nblcks, nthrds, 0, strm>>>
698 ((real *) u1, (real *) u2, (real *) u3,
699 (real *) v1, (real *) v2, (real *) v3,
700 (real *) w1, (real *) w2, (real *) w3, *n);
701 CUDA_CHECK(cudaGetLastError());
702 }
703
704
705 /*
706 * Reduction buffers, owned by the device layer and released
707 * on device teardown (cuda_buffer_free_all in cuda_finalize)
708 */
709 cuda_buffer_t redbuf = CUDA_BUFFER_INIT_SYMM;
710 cuda_buffer_t redbuf_xp = CUDA_BUFFER_INIT_SYMM;
711
715 void cuda_redbuf_check_alloc(int nb) {
716 cuda_buffer_reserve(&redbuf, (nb + 1) * sizeof(real));
717 }
718
722 void cuda_redbuf_check_alloc_xp(int nb) {
723 cuda_buffer_reserve(&redbuf_xp, (nb + 1) * sizeof(real_xp));
724 }
725
729 void cuda_global_reduce_add(real * bufred, void * bufred_d,
730 int n, const cudaStream_t stream) {
731
732#ifdef HAVE_NCCL
733 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
734 DEVICE_NCCL_SUM, stream);
735 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real)*n,
736 cudaMemcpyDeviceToHost, stream));
737 cudaStreamSynchronize(stream);
738#elif HAVE_NVSHMEM
739 if (sizeof(real) == sizeof(float)) {
740 nvshmemx_float_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
741 (float *) bufred_d,
742 (float *) bufred_d, n, stream);
743 }
744 else if (sizeof(real) == sizeof(double)) {
745 nvshmemx_double_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
746 (double *) bufred_d,
747 (double *) bufred_d, n, stream);
748
749 }
750 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
751 sizeof(real)*n, cudaMemcpyDeviceToHost, stream));
752 cudaStreamSynchronize(stream);
753#elif HAVE_DEVICE_MPI
754 cudaStreamSynchronize(stream);
755 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_SUM);
756#else
757 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real),
758 cudaMemcpyDeviceToHost, stream));
759 cudaStreamSynchronize(stream);
760#endif
761 }
762
766 void cuda_global_reduce_add_xp(real_xp * bufred, void * bufred_d,
767 int n, const cudaStream_t stream) {
768
769#ifdef HAVE_NCCL
770 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real_xp),
771 DEVICE_NCCL_SUM, stream);
772 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real_xp)*n,
773 cudaMemcpyDeviceToHost, stream));
774 cudaStreamSynchronize(stream);
775#elif HAVE_NVSHMEM
776 if (sizeof(real_xp) == sizeof(float)) {
777 nvshmemx_float_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
778 (float *) bufred_d,
779 (float *) bufred_d, n, stream);
780 }
781 else if (sizeof(real_xp) == sizeof(double)) {
782 nvshmemx_double_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
783 (double *) bufred_d,
784 (double *) bufred_d, n, stream);
785
786 }
787 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
788 sizeof(real_xp)*n, cudaMemcpyDeviceToHost, stream));
789 cudaStreamSynchronize(stream);
790#elif HAVE_DEVICE_MPI
791 cudaStreamSynchronize(stream);
792 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real_xp), DEVICE_MPI_SUM);
793#else
794 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real_xp),
795 cudaMemcpyDeviceToHost, stream));
796 cudaStreamSynchronize(stream);
797#endif
798 }
799
803 void cuda_global_reduce_max(real * bufred, void * bufred_d,
804 int n, const cudaStream_t stream) {
805
806#ifdef HAVE_NCCL
807 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
808 DEVICE_NCCL_MAX, stream);
809 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real)*n,
810 cudaMemcpyDeviceToHost, stream));
811 cudaStreamSynchronize(stream);
812#elif HAVE_NVSHMEM
813 if (sizeof(real) == sizeof(float)) {
814 nvshmemx_float_max_reduce_on_stream(NVSHMEM_TEAM_WORLD,
815 (float *) bufred_d,
816 (float *) bufred_d, n, stream);
817 }
818 else if (sizeof(real) == sizeof(double)) {
819 nvshmemx_double_max_reduce_on_stream(NVSHMEM_TEAM_WORLD,
820 (double *) bufred_d,
821 (double *) bufred_d, n, stream);
822
823 }
824 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
825 sizeof(real)*n, cudaMemcpyDeviceToHost, stream));
826 cudaStreamSynchronize(stream);
827#elif HAVE_DEVICE_MPI
828 cudaStreamSynchronize(stream);
829 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MAX);
830#else
831 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real),
832 cudaMemcpyDeviceToHost, stream));
833 cudaStreamSynchronize(stream);
834#endif
835 }
836
840 void cuda_global_reduce_min(real * bufred, void * bufred_d,
841 int n, const cudaStream_t stream) {
842
843#ifdef HAVE_NCCL
844 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
845 DEVICE_NCCL_MIN, stream);
846 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real)*n,
847 cudaMemcpyDeviceToHost, stream));
848 cudaStreamSynchronize(stream);
849#elif HAVE_NVSHMEM
850 if (sizeof(real) == sizeof(float)) {
851 nvshmemx_float_min_reduce_on_stream(NVSHMEM_TEAM_WORLD,
852 (float *) bufred_d,
853 (float *) bufred_d, n, stream);
854 }
855 else if (sizeof(real) == sizeof(double)) {
856 nvshmemx_double_min_reduce_on_stream(NVSHMEM_TEAM_WORLD,
857 (double *) bufred_d,
858 (double *) bufred_d, n, stream);
859
860 }
861 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
862 sizeof(real)*n, cudaMemcpyDeviceToHost, stream));
863 cudaStreamSynchronize(stream);
864#elif HAVE_DEVICE_MPI
865 cudaStreamSynchronize(stream);
866 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MIN);
867#else
868 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real),
869 cudaMemcpyDeviceToHost, stream));
870 cudaStreamSynchronize(stream);
871#endif
872 }
873
874
879 real cuda_vlsc3(void *u, void *v, void *w, int *n, cudaStream_t stream) {
880
881 const dim3 nthrds(1024, 1, 1);
882 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
883 const int nb = ((*n) + 1024 - 1)/ 1024;
884
885 cuda_redbuf_check_alloc(nb);
886
887 vlsc3_kernel<real><<<nblcks, nthrds, 0, stream>>>
888 ((real *) u, (real *) v, (real *) w, (real *) redbuf.dev, *n);
889 CUDA_CHECK(cudaGetLastError());
890 reduce_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev, nb);
891 CUDA_CHECK(cudaGetLastError());
892
893 CUDA_CHECK(cudaMemcpyAsync(((real *) redbuf.host), redbuf.dev, sizeof(real),
894 cudaMemcpyDeviceToHost, stream));
895 cudaStreamSynchronize(stream);
896
897 return ((real *) redbuf.host)[0];
898 }
899
900
901
902
907 real_xp cuda_glsc3(void *a, void *b, void *c, int *n, cudaStream_t stream) {
908
909 const dim3 nthrds(1024, 1, 1);
910 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
911 const int nb = ((*n) + 1024 - 1)/ 1024;
912
913 cuda_redbuf_check_alloc_xp(nb);
914
915 if ( *n > 0) {
916 glsc3_kernel<real, real_xp><<<nblcks, nthrds, 0, stream>>>
917 ((real *) a, (real *) b, (real *) c, (real_xp *) redbuf_xp.dev, *n);
918 CUDA_CHECK(cudaGetLastError());
919 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
920 CUDA_CHECK(cudaGetLastError());
921 }
922 else {
923 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
924 }
925 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
926
927 return ((real_xp *) redbuf_xp.host)[0];
928 }
929
934 void cuda_glsc3_many(real_xp *h, void * w, void *v,void *mult, int *j, int *n,
935 cudaStream_t stream){
936 int pow2 = 1;
937 while(pow2 < (*j)){
938 pow2 = 2*pow2;
939 }
940 const int nt = 1024/pow2;
941 const dim3 nthrds(pow2, nt, 1);
942 const dim3 nblcks(((*n)+nt - 1)/nt, 1, 1);
943 const int nb = ((*n) + nt - 1)/nt;
944
945 cuda_redbuf_check_alloc_xp((*j)*nb);
946
947 if ( *n > 0) {
948 glsc3_many_kernel<real, real_xp><<<nblcks, nthrds, 0, stream>>>
949 ((const real *) w, (const real **) v,
950 (const real *)mult, (real_xp *)redbuf_xp.dev, *j, *n);
951 CUDA_CHECK(cudaGetLastError());
952 glsc3_reduce_kernel<real_xp>
953 <<<(*j), 1024, 0, stream>>>((real_xp *) redbuf_xp.dev, nb, *j);
954 CUDA_CHECK(cudaGetLastError());
955 }
956 else {
957 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
958 }
959 cuda_global_reduce_add_xp(h, redbuf_xp.dev, (*j), stream);
960 }
961
966 real_xp cuda_glsc2(void *a, void *b, int *n, cudaStream_t stream) {
967
968 const dim3 nthrds(1024, 1, 1);
969 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
970 const int nb = ((*n) + 1024 - 1)/ 1024;
971
972 cuda_redbuf_check_alloc_xp(nb);
973
974 if ( *n > 0) {
975 glsc2_kernel<real, real_xp>
976 <<<nblcks, nthrds, 0, stream>>>((real *) a,
977 (real *) b,
978 (real_xp *) redbuf_xp.dev, *n);
979 CUDA_CHECK(cudaGetLastError());
980 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
981 CUDA_CHECK(cudaGetLastError());
982 }
983 else {
984 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
985 }
986 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
987
988 return ((real_xp *) redbuf_xp.host)[0];
989 }
990
995 real_xp cuda_glsubnorm2(void *a, void *b, int *n, cudaStream_t stream) {
996
997 const dim3 nthrds(1024, 1, 1);
998 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
999 const int nb = ((*n) + 1024 - 1)/ 1024;
1000
1001 cuda_redbuf_check_alloc_xp(nb);
1002
1003 if ( *n > 0) {
1004 glsubnorm2_kernel<real, real_xp>
1005 <<<nblcks, nthrds, 0, stream>>>((real *) a,
1006 (real *) b,
1007 (real_xp *) redbuf_xp.dev, *n);
1008 CUDA_CHECK(cudaGetLastError());
1009 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
1010 CUDA_CHECK(cudaGetLastError());
1011 }
1012 else {
1013 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1014 }
1015 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
1016
1017 return ((real_xp *) redbuf_xp.host)[0];
1018 }
1019
1024 real_xp cuda_glsum(void *a, int *n, cudaStream_t stream) {
1025 const dim3 nthrds(1024, 1, 1);
1026 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1027 const int nb = ((*n) + 1024 - 1)/ 1024;
1028
1029 cuda_redbuf_check_alloc_xp(nb);
1030 if ( *n > 0) {
1031 glsum_kernel<real, real_xp>
1032 <<<nblcks, nthrds, 0, stream>>>((real *) a,
1033 (real_xp *) redbuf_xp.dev, *n);
1034 CUDA_CHECK(cudaGetLastError());
1035 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
1036 CUDA_CHECK(cudaGetLastError());
1037 }
1038 else {
1039 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1040 }
1041
1042 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
1043
1044 return ((real_xp *) redbuf_xp.host)[0];
1045 }
1046
1051 real cuda_glmax(void *a, real *ninf, int *n, cudaStream_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 cuda_redbuf_check_alloc(nb);
1057 if ( *n > 0) {
1058 glmax_kernel<real>
1059 <<<nblcks, nthrds, 0, stream>>>((real *) a, *ninf,
1060 (real *) redbuf.dev, *n);
1061 CUDA_CHECK(cudaGetLastError());
1062 reduce_max_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev,
1063 *ninf, nb);
1064 CUDA_CHECK(cudaGetLastError());
1065 }
1066 else {
1067 int nel = (int) (redbuf.size / sizeof(real));
1068 cuda_rzero(redbuf.dev, &nel, stream);
1069 }
1070
1071 cuda_global_reduce_max(((real *) redbuf.host), redbuf.dev, 1, stream);
1072
1073 return ((real *) redbuf.host)[0];
1074 }
1075
1080 real cuda_glmin(void *a, real *pinf, int *n, cudaStream_t stream) {
1081 const dim3 nthrds(1024, 1, 1);
1082 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1083 const int nb = ((*n) + 1024 - 1)/ 1024;
1084
1085 cuda_redbuf_check_alloc(nb);
1086 if ( *n > 0) {
1087 glmin_kernel<real>
1088 <<<nblcks, nthrds, 0, stream>>>((real *) a, *pinf,
1089 (real *) redbuf.dev, *n);
1090 CUDA_CHECK(cudaGetLastError());
1091 reduce_min_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev,
1092 *pinf, nb);
1093 CUDA_CHECK(cudaGetLastError());
1094 }
1095 else {
1096 int nel = (int) (redbuf.size / sizeof(real));
1097 cuda_rzero(redbuf.dev, &nel, stream);
1098 }
1099
1100 cuda_global_reduce_min(((real *) redbuf.host), redbuf.dev, 1, stream);
1101
1102 return ((real *) redbuf.host)[0];
1103 }
1104
1109 real cuda_glamax(void *a, int *n, cudaStream_t stream) {
1110 const dim3 nthrds(1024, 1, 1);
1111 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1112 const int nb = ((*n) + 1024 - 1)/ 1024;
1113 const real zero = 0.0;
1114
1115 cuda_redbuf_check_alloc(nb);
1116 if ( *n > 0) {
1117 glamax_kernel<real>
1118 <<<nblcks, nthrds, 0, stream>>>((real *) a,
1119 (real *) redbuf.dev, *n);
1120 CUDA_CHECK(cudaGetLastError());
1121 reduce_max_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev,
1122 zero, nb);
1123 CUDA_CHECK(cudaGetLastError());
1124 }
1125 else {
1126 int nel = (int) (redbuf.size / sizeof(real));
1127 cuda_rzero(redbuf.dev, &nel, stream);
1128 }
1129
1130 cuda_global_reduce_max(((real *) redbuf.host), redbuf.dev, 1, stream);
1131
1132 return ((real *) redbuf.host)[0];
1133 }
1134
1139 void cuda_absval(void *a, int *n, cudaStream_t stream) {
1140
1141 const dim3 nthrds(1024, 1, 1);
1142 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1143
1144 absval_kernel<real>
1145 <<<nblcks, nthrds,0, stream>>>((real *) a, * n);
1146 CUDA_CHECK(cudaGetLastError());
1147
1148 }
1149
1150 // ======================================================================== //
1151 // Point-wise operations.
1152
1157 void cuda_pwmax_vec2(void *a, void *b, int *n, cudaStream_t stream) {
1158
1159 const dim3 nthrds(1024, 1, 1);
1160 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1161
1162 pwmax_vec2_kernel<real><<<nblcks, nthrds, 0, stream>>>
1163 ((real *)a, (real *)b, *n);
1164 CUDA_CHECK(cudaGetLastError());
1165 }
1166
1171 void cuda_pwmax_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream) {
1172
1173 const dim3 nthrds(1024, 1, 1);
1174 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1175
1176 pwmax_vec3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1177 ((real *)a, (real *)b, (real *)c, *n);
1178 CUDA_CHECK(cudaGetLastError());
1179 }
1180
1185 void cuda_pwmax_sca2(void *a, real *c, int *n, cudaStream_t stream) {
1186
1187 const dim3 nthrds(1024, 1, 1);
1188 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1189
1190 pwmax_sca2_kernel<real><<<nblcks, nthrds, 0, stream>>>
1191 ((real *)a, *c, *n);
1192 CUDA_CHECK(cudaGetLastError());
1193 }
1194
1199 void cuda_pwmax_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream) {
1200
1201 const dim3 nthrds(1024, 1, 1);
1202 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1203
1204 pwmax_sca3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1205 ((real *)a, (real *)b, *c, *n);
1206 CUDA_CHECK(cudaGetLastError());
1207 }
1208
1213 void cuda_pwmin_vec2(void *a, void *b, int *n, cudaStream_t stream) {
1214
1215 const dim3 nthrds(1024, 1, 1);
1216 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1217
1218 pwmin_vec2_kernel<real><<<nblcks, nthrds, 0, stream>>>
1219 ((real *)a, (real *)b, *n);
1220 CUDA_CHECK(cudaGetLastError());
1221 }
1222
1227 void cuda_pwmin_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream) {
1228
1229 const dim3 nthrds(1024, 1, 1);
1230 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1231
1232 pwmin_vec3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1233 ((real *)a, (real *)b, (real *)c, *n);
1234 CUDA_CHECK(cudaGetLastError());
1235 }
1236
1241 void cuda_pwmin_sca2(void *a, real *c, int *n, cudaStream_t stream) {
1242
1243 const dim3 nthrds(1024, 1, 1);
1244 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1245
1246 pwmin_sca2_kernel<real><<<nblcks, nthrds, 0, stream>>>((real *)a, *c, *n);
1247 CUDA_CHECK(cudaGetLastError());
1248 }
1249
1254 void cuda_pwmin_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream) {
1255
1256 const dim3 nthrds(1024, 1, 1);
1257 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1258
1259 pwmin_sca3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1260 ((real *)a, (real *)b, *c, *n);
1261 CUDA_CHECK(cudaGetLastError());
1262 }
1263
1264 // ======================================================================== //
1265
1269 void cuda_iadd(void *a, int *c, int *n, cudaStream_t stream) {
1270
1271 const dim3 nthrds(1024, 1, 1);
1272 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1273
1274 cadd_kernel<int><<<nblcks, nthrds, 0, stream>>>
1275 ((int *) a, *c, *n);
1276 CUDA_CHECK(cudaGetLastError());
1277
1278 }
1279
1280} /* extern "C" */
cuda_buffer_reserve
void cuda_buffer_reserve(cuda_buffer_t *buf, size_t size)
Definition buffer.cu:50
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
CUDA_BUFFER_INIT_SYMM
#define CUDA_BUFFER_INIT_SYMM
Definition buffer.h:63
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
CUDA_CHECK
#define CUDA_CHECK(err)
Definition check.h:6
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
math_kernel.h
cuda_global_reduce_min
void cuda_global_reduce_min(real *bufred, void *bufred_d, int n, const cudaStream_t stream)
Definition math.cu:840
cuda_pwmax_sca3
void cuda_pwmax_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream)
Definition math.cu:1199
cuda_sqrt_inplace
void cuda_sqrt_inplace(void *a, int *n, cudaStream_t strm)
Definition math.cu:311
cuda_addcol4
void cuda_addcol4(void *a, void *b, void *c, void *d, int *n, cudaStream_t strm)
Definition math.cu:642
cuda_add2s2
void cuda_add2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm)
Definition math.cu:416
cuda_vdot3
void cuda_vdot3(void *dot, void *u1, void *u2, void *u3, void *v1, void *v2, void *v3, int *n, cudaStream_t strm)
Definition math.cu:672
cuda_add3s2
void cuda_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n, cudaStream_t strm)
Definition math.cu:466
cuda_face_masked_gather_copy
void cuda_face_masked_gather_copy(void *a, void *b, void *mask, void *facet, int *n1, int *n2, int *lx, int *ly, int *lz, int *m, cudaStream_t strm)
Definition math.cu:130
cuda_global_reduce_add
void cuda_global_reduce_add(real *bufred, void *bufred_d, int n, const cudaStream_t stream)
Definition math.cu:729
cuda_addcol3
void cuda_addcol3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:628
cuda_pwmin_vec2
void cuda_pwmin_vec2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:1213
cuda_pwmax_vec3
void cuda_pwmax_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream)
Definition math.cu:1171
cuda_power
void cuda_power(void *ap, void *a, real *p, int *n, cudaStream_t strm)
Definition math.cu:325
cuda_glmin
real cuda_glmin(void *a, real *pinf, int *n, cudaStream_t stream)
Definition math.cu:1080
cuda_cmult2
void cuda_cmult2(void *a, void *b, real *c, int *n, cudaStream_t strm)
Definition math.cu:226
cuda_add2s1
void cuda_add2s1(void *a, void *b, real *c1, int *n, cudaStream_t strm)
Definition math.cu:400
cuda_col3
void cuda_col3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:571
cuda_masked_copy_0
void cuda_masked_copy_0(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:69
cuda_add2s2_many
void cuda_add2s2_many(void *x, void **p, void *alpha, int *j, int *n, cudaStream_t strm)
Definition math.cu:433
cuda_sub3
void cuda_sub3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:614
cuda_copy
void cuda_copy(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:60
cuda_cfill_mask
void cuda_cfill_mask(void *a, real *c, int *size, int *mask, int *mask_size, cudaStream_t strm)
Definition math.cu:192
cuda_glsc3
real_xp cuda_glsc3(void *a, void *b, void *c, int *n, cudaStream_t stream)
Definition math.cu:907
cuda_masked_copy_aligned
void cuda_masked_copy_aligned(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:85
cuda_invcol2
void cuda_invcol2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:530
cuda_global_reduce_add_xp
void cuda_global_reduce_add_xp(real_xp *bufred, void *bufred_d, int n, const cudaStream_t stream)
Definition math.cu:766
cuda_pwmax_sca2
void cuda_pwmax_sca2(void *a, real *c, int *n, cudaStream_t stream)
Definition math.cu:1185
cuda_col2
void cuda_col2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:558
cuda_masked_scatter_copy
void cuda_masked_scatter_copy(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:163
cuda_add2
void cuda_add2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:355
cuda_redbuf_check_alloc
void cuda_redbuf_check_alloc(int nb)
Definition math.cu:715
cuda_cmult
void cuda_cmult(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:213
cuda_add4
void cuda_add4(void *a, void *b, void *c, void *d, int *n, cudaStream_t strm)
Definition math.cu:384
cuda_masked_scatter_copy_aligned
void cuda_masked_scatter_copy_aligned(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:177
cuda_sub2
void cuda_sub2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:600
cuda_redbuf_check_alloc_xp
void cuda_redbuf_check_alloc_xp(int nb)
Definition math.cu:722
cuda_invcol3
void cuda_invcol3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:544
cuda_vcross
void cuda_vcross(void *u1, void *u2, void *u3, void *v1, void *v2, void *v3, void *w1, void *w2, void *w3, int *n, cudaStream_t strm)
Definition math.cu:689
cuda_glsc2
real_xp cuda_glsc2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:966
cuda_masked_gather_copy
void cuda_masked_gather_copy(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:100
cuda_vlsc3
real cuda_vlsc3(void *u, void *v, void *w, int *n, cudaStream_t stream)
Definition math.cu:879
cuda_pwmin_sca3
void cuda_pwmin_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream)
Definition math.cu:1254
cuda_cadd2
void cuda_cadd2(void *a, void *b, real *c, int *n, cudaStream_t strm)
Definition math.cu:281
cuda_addsqr2s2
void cuda_addsqr2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm)
Definition math.cu:450
redbuf_xp
cuda_buffer_t redbuf_xp
Definition math.cu:710
cuda_glsum
real_xp cuda_glsum(void *a, int *n, cudaStream_t stream)
Definition math.cu:1024
cuda_add3
void cuda_add3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:370
cuda_addcol3s2
void cuda_addcol3s2(void *a, void *b, void *c, real *s, int *n, cudaStream_t strm)
Definition math.cu:657
cuda_cwrap
void cuda_cwrap(void *a, real *min_val, real *max_val, int *n, cudaStream_t strm)
Definition math.cu:296
cuda_pwmin_sca2
void cuda_pwmin_sca2(void *a, real *c, int *n, cudaStream_t stream)
Definition math.cu:1241
cuda_cfill
void cuda_cfill(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:339
redbuf
cuda_buffer_t redbuf
Definition math.cu:709
cuda_absval
void cuda_absval(void *a, int *n, cudaStream_t stream)
Definition math.cu:1139
cuda_rzero
void cuda_rzero(void *a, int *n, cudaStream_t strm)
Definition math.cu:206
cuda_cdiv2
void cuda_cdiv2(void *a, void *b, real *c, int *n, cudaStream_t strm)
Definition math.cu:253
cuda_cdiv
void cuda_cdiv(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:240
cuda_glamax
real cuda_glamax(void *a, int *n, cudaStream_t stream)
Definition math.cu:1109
cuda_global_reduce_max
void cuda_global_reduce_max(real *bufred, void *bufred_d, int n, const cudaStream_t stream)
Definition math.cu:803
cuda_iadd
void cuda_iadd(void *a, int *c, int *n, cudaStream_t stream)
Definition math.cu:1269
cuda_masked_gather_copy_aligned
void cuda_masked_gather_copy_aligned(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:115
cuda_pwmax_vec2
void cuda_pwmax_vec2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:1157
cuda_glsc3_many
void cuda_glsc3_many(real_xp *h, void *w, void *v, void *mult, int *j, int *n, cudaStream_t stream)
Definition math.cu:934
cuda_glsubnorm2
real_xp cuda_glsubnorm2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:995
cuda_masked_atomic_reduction
void cuda_masked_atomic_reduction(void *a, void *b, void *mask, int *n, int *m, cudaStream_t strm)
Definition math.cu:148
cuda_add5s4
void cuda_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1, real *c2, real *c3, real *c4, int *n, cudaStream_t strm)
Definition math.cu:500
cuda_invcol1
void cuda_invcol1(void *a, int *n, cudaStream_t strm)
Definition math.cu:517
cuda_radd
void cuda_radd(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:267
cuda_add4s3
void cuda_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2, real *c3, int *n, cudaStream_t strm)
Definition math.cu:483
cuda_glmax
real cuda_glmax(void *a, real *ninf, int *n, cudaStream_t stream)
Definition math.cu:1051
cuda_subcol3
void cuda_subcol3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:585
cuda_pwmin_vec3
void cuda_pwmin_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream)
Definition math.cu:1227
mask
Object for handling masks in Neko.
Definition mask.f90:34
cuda_buffer
Definition buffer.h:51
cuda_buffer::host
void * host
Definition buffer.h:52
cuda_buffer::size
size_t size
Definition buffer.h:54
cuda_buffer::dev
void * dev
Definition buffer.h:53