Neko 1.99.6
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-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 "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(1024, 1, 1);
152 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
153
154 masked_atomic_reduction_kernel<real><<<nblcks, nthrds, 0, strm>>>
155 ((real *) a, (real *) b, (int *) mask, *n, *m);
156 CUDA_CHECK(cudaGetLastError());
157
158 }
162 void cuda_masked_scatter_copy(void *a, void *b, void *mask,
163 int *n, int *m, cudaStream_t strm) {
164
165 const dim3 nthrds(1024, 1, 1);
166 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
167
168 masked_scatter_copy_kernel<real><<<nblcks, nthrds, 0, strm>>>
169 ((real *) a, (real*) b,(int*) mask, *n, *m);
170 CUDA_CHECK(cudaGetLastError());
171 }
172
176 void cuda_masked_scatter_copy_aligned(void *a, void *b, void *mask,
177 int *n, int *m, cudaStream_t strm) {
178
179 const dim3 nthrds(1024, 1, 1);
180 const dim3 nblcks(((*m)+1024 - 1)/ 1024, 1, 1);
181
182 masked_scatter_copy_aligned_kernel<real><<<nblcks, nthrds, 0, strm>>>
183 ((real *) a, (real*) b,(int*) mask, *n, *m);
184 CUDA_CHECK(cudaGetLastError());
185 }
186
187
191 void cuda_cfill_mask(void* a, real* c, int* size, int* mask, int* mask_size,
192 cudaStream_t strm) {
193
194 const dim3 nthrds(1024, 1, 1);
195 const dim3 nblcks(((*mask_size) + 1024 - 1) / 1024, 1, 1);
196
197 cfill_mask_kernel<real><<<nblcks, nthrds, 0, strm>>>
198 ((real*)a, *c, *size, mask, *mask_size);
199 CUDA_CHECK(cudaGetLastError());
200 }
201
205 void cuda_rzero(void *a, int *n, cudaStream_t strm) {
206 CUDA_CHECK(cudaMemsetAsync(a, 0, (*n) * sizeof(real), strm));
207 }
208
212 void cuda_cmult(void *a, real *c, int *n, cudaStream_t strm) {
213
214 const dim3 nthrds(1024, 1, 1);
215 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
216
217 cmult_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
218 CUDA_CHECK(cudaGetLastError());
219
220 }
221
225 void cuda_cmult2(void *a, void *b, real *c, int *n, cudaStream_t strm) {
226
227 const dim3 nthrds(1024, 1, 1);
228 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
229
230 cmult2_kernel<real><<<nblcks, nthrds, 0, strm>>>
231 ((real *) a, (real *) b, *c, *n);
232 CUDA_CHECK(cudaGetLastError());
233
234 }
235
239 void cuda_cdiv(void *a, real *c, int *n, cudaStream_t strm) {
240
241 const dim3 nthrds(1024, 1, 1);
242 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
243
244 cdiv_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
245 CUDA_CHECK(cudaGetLastError());
246
247 }
248
252 void cuda_cdiv2(void *a, void *b, real *c, int *n, cudaStream_t strm) {
253
254 const dim3 nthrds(1024, 1, 1);
255 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
256
257 cdiv2_kernel<real><<<nblcks, nthrds, 0, strm>>>
258 ((real *) a, (real *) b, *c, *n);
259 CUDA_CHECK(cudaGetLastError());
260
261 }
262
266 void cuda_radd(void *a, real *c, int *n, cudaStream_t strm) {
267
268 const dim3 nthrds(1024, 1, 1);
269 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
270
271 cadd_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
272 CUDA_CHECK(cudaGetLastError());
273
274 }
275
280 void cuda_cadd2(void *a, void *b, real *c, int *n, cudaStream_t strm) {
281
282 const dim3 nthrds(1024, 1, 1);
283 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
284
285 cadd2_kernel<real><<<nblcks, nthrds, 0, strm>>>
286 ((real *) a, (real *) b, *c, *n);
287 CUDA_CHECK(cudaGetLastError());
288
289 }
290
295 void cuda_cwrap(void *a, real *min_val, real *max_val, int *n,
296 cudaStream_t strm) {
297
298 const dim3 nthrds(1024, 1, 1);
299 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
300
301 cwrap_kernel<real><<<nblcks, nthrds, 0, strm>>>
302 ((real *) a, *min_val, *max_val, *n);
303 CUDA_CHECK(cudaGetLastError());
304
305 }
306
310 void cuda_sqrt_inplace(void *a, int *n, cudaStream_t strm) {
311
312 const dim3 nthrds(1024, 1, 1);
313 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
314
315 sqrt_inplace_kernel<real><<<nblcks, nthrds, 0, strm>>>
316 ((real *) a, *n);
317 CUDA_CHECK(cudaGetLastError());
318
319 }
320
324 void cuda_power(void *ap, void *a, real *p, int *n, cudaStream_t strm) {
325
326 const dim3 nthrds(1024, 1, 1);
327 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
328
329 power_kernel<real><<<nblcks, nthrds, 0, strm>>>
330 ((real *) ap, (real *) a, *p, *n);
331 CUDA_CHECK(cudaGetLastError());
332
333 }
334
338 void cuda_cfill(void *a, real *c, int *n, cudaStream_t strm) {
339
340 const dim3 nthrds(1024, 1, 1);
341 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
342
343 if (*n > 0){
344 cfill_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *c, *n);
345 CUDA_CHECK(cudaGetLastError());
346 }
347
348 }
349
354 void cuda_add2(void *a, void *b, int *n, cudaStream_t strm) {
355
356 const dim3 nthrds(1024, 1, 1);
357 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
358
359 add2_kernel<real><<<nblcks, nthrds, 0, strm>>>
360 ((real *) a, (real *) b, *n);
361 CUDA_CHECK(cudaGetLastError());
362
363 }
364
369 void cuda_add3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
370
371 const dim3 nthrds(1024, 1, 1);
372 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
373
374 add3_kernel<real><<<nblcks, nthrds, 0, strm>>>
375 ((real *) a, (real *) b, (real *) c, *n);
376 CUDA_CHECK(cudaGetLastError());
377 }
378
383 void cuda_add4(void *a, void *b, void *c, void *d, int *n,
384 cudaStream_t strm) {
385
386 const dim3 nthrds(1024, 1, 1);
387 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
388
389 add4_kernel<real><<<nblcks, nthrds, 0, strm>>>
390 ((real *) a, (real *) b, (real *) c, (real *) d, *n);
391 CUDA_CHECK(cudaGetLastError());
392
393 }
399 void cuda_add2s1(void *a, void *b, real *c1, int *n, cudaStream_t strm) {
400
401 const dim3 nthrds(1024, 1, 1);
402 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
403
404 add2s1_kernel<real><<<nblcks, nthrds, 0, strm>>>
405 ((real *) a, (real *) b, *c1, *n);
406 CUDA_CHECK(cudaGetLastError());
407
408 }
409
415 void cuda_add2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm) {
416
417 const dim3 nthrds(1024, 1, 1);
418 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
419
420 add2s2_kernel<real><<<nblcks, nthrds, 0, strm>>>
421 ((real *) a, (real *) b, *c1, *n);
422 CUDA_CHECK(cudaGetLastError());
423
424 }
425
432 void cuda_add2s2_many(void *x, void **p, void *alpha, int *j, int *n,
433 cudaStream_t strm) {
434
435 const dim3 nthrds(1024, 1, 1);
436 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
437
438 add2s2_many_kernel<real><<<nblcks, nthrds, 0, strm>>>
439 ((real *) x, (const real **) p, (real *) alpha, *j, *n);
440 CUDA_CHECK(cudaGetLastError());
441
442 }
443
449 void cuda_addsqr2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm) {
450
451 const dim3 nthrds(1024, 1, 1);
452 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
453
454 addsqr2s2_kernel<real><<<nblcks, nthrds, 0,strm>>>
455 ((real *) a, (real *) b, *c1, *n);
456 CUDA_CHECK(cudaGetLastError());
457
458 }
459
465 void cuda_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n,
466 cudaStream_t strm) {
467
468 const dim3 nthrds(1024, 1, 1);
469 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
470
471 add3s2_kernel<real><<<nblcks, nthrds, 0, strm>>>
472 ((real *) a, (real *) b, (real *) c, *c1, *c2, *n);
473 CUDA_CHECK(cudaGetLastError());
474
475 }
476
482 void cuda_add4s3(void *a, void *b, void *c, void *d, real *c1, real *c2,
483 real *c3, int *n, cudaStream_t strm) {
484
485 const dim3 nthrds(1024, 1, 1);
486 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
487
488 add4s3_kernel<real><<<nblcks, nthrds, 0, strm>>>
489 ((real *) a, (real *) b, (real *) c, (real *) d, *c1, *c2, *c3, *n);
490 CUDA_CHECK(cudaGetLastError());
491
492 }
493
499 void cuda_add5s4(void *a, void *b, void *c, void *d, void *e, real *c1,
500 real *c2, real *c3, real *c4, int *n, cudaStream_t strm) {
501
502 const dim3 nthrds(1024, 1, 1);
503 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
504
505 add5s4_kernel<real><<<nblcks, nthrds, 0, strm>>>
506 ((real *) a, (real *) b, (real *) c, (real *) d, (real *) e,
507 *c1, *c2, *c3, *c4, *n);
508 CUDA_CHECK(cudaGetLastError());
509
510 }
511
516 void cuda_invcol1(void *a, int *n, cudaStream_t strm) {
517
518 const dim3 nthrds(1024, 1, 1);
519 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
520
521 invcol1_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, *n);
522 CUDA_CHECK(cudaGetLastError());
523 }
524
529 void cuda_invcol2(void *a, void *b, int *n, cudaStream_t strm) {
530
531 const dim3 nthrds(1024, 1, 1);
532 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
533
534 invcol2_kernel<real><<<nblcks, nthrds, 0, strm>>>
535 ((real *) a, (real *) b, *n);
536 CUDA_CHECK(cudaGetLastError());
537 }
538
543 void cuda_invcol3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
544
545 const dim3 nthrds(1024, 1, 1);
546 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
547
548 invcol3_kernel<real><<<nblcks, nthrds, 0, strm>>>
549 ((real *) a, (real *) b, (real *) c, *n);
550 CUDA_CHECK(cudaGetLastError());
551 }
552
557 void cuda_col2(void *a, void *b, int *n, cudaStream_t strm) {
558
559 const dim3 nthrds(1024, 1, 1);
560 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
561
562 col2_kernel<real><<<nblcks, nthrds, 0, strm>>>((real *) a, (real *) b, *n);
563 CUDA_CHECK(cudaGetLastError());
564 }
565
570 void cuda_col3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
571
572 const dim3 nthrds(1024, 1, 1);
573 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
574
575 col3_kernel<real><<<nblcks, nthrds, 0, strm>>>
576 ((real *) a, (real *) b, (real *) c, *n);
577 CUDA_CHECK(cudaGetLastError());
578 }
579
584 void cuda_subcol3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
585
586 const dim3 nthrds(1024, 1, 1);
587 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
588
589 subcol3_kernel<real><<<nblcks, nthrds, 0, strm>>>
590 ((real *) a, (real *) b, (real *) c, *n);
591 CUDA_CHECK(cudaGetLastError());
592 }
593
594
599 void cuda_sub2(void *a, void *b, int *n, cudaStream_t strm) {
600
601 const dim3 nthrds(1024, 1, 1);
602 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
603
604 sub2_kernel<real><<<nblcks, nthrds, 0, strm>>>
605 ((real *) a, (real *) b, *n);
606 CUDA_CHECK(cudaGetLastError());
607 }
608
613 void cuda_sub3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
614
615 const dim3 nthrds(1024, 1, 1);
616 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
617
618 sub3_kernel<real><<<nblcks, nthrds, 0, strm>>>
619 ((real *) a, (real *) b, (real *) c, *n);
620 CUDA_CHECK(cudaGetLastError());
621 }
622
627 void cuda_addcol3(void *a, void *b, void *c, int *n, cudaStream_t strm) {
628
629 const dim3 nthrds(1024, 1, 1);
630 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
631
632 addcol3_kernel<real><<<nblcks, nthrds, 0, strm>>>
633 ((real *) a, (real *) b, (real *) c, *n);
634 CUDA_CHECK(cudaGetLastError());
635 }
636
641 void cuda_addcol4(void *a, void *b, void *c, void *d, int *n,
642 cudaStream_t strm) {
643
644 const dim3 nthrds(1024, 1, 1);
645 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
646
647 addcol4_kernel<real><<<nblcks, nthrds, 0, strm>>>
648 ((real *) a, (real *) b, (real *) c, (real *) d, *n);
649 CUDA_CHECK(cudaGetLastError());
650 }
651
656 void cuda_addcol3s2(void *a, void *b, void *c, real *s, int *n,
657 cudaStream_t strm) {
658
659 const dim3 nthrds(1024, 1, 1);
660 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
661
662 addcol3s2_kernel<real><<<nblcks, nthrds, 0, strm>>>
663 ((real *) a, (real *) b, (real *) c, *s, *n);
664 CUDA_CHECK(cudaGetLastError());
665 }
666
671 void cuda_vdot3(void *dot, void *u1, void *u2, void *u3,
672 void *v1, void *v2, void *v3, int *n,
673 cudaStream_t strm) {
674
675 const dim3 nthrds(1024, 1, 1);
676 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
677
678 vdot3_kernel<real><<<nblcks, nthrds, 0, strm>>>
679 ((real *) dot, (real *) u1, (real *) u2, (real *) u3,
680 (real *) v1, (real *) v2, (real *) v3, *n);
681 CUDA_CHECK(cudaGetLastError());
682 }
683
688 void cuda_vcross(void *u1, void *u2, void *u3,
689 void *v1, void *v2, void *v3,
690 void *w1, void *w2, void *w3,
691 int *n, cudaStream_t strm) {
692
693 const dim3 nthrds(1024, 1, 1);
694 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
695
696 vcross_kernel<real><<<nblcks, nthrds, 0, strm>>>
697 ((real *) u1, (real *) u2, (real *) u3,
698 (real *) v1, (real *) v2, (real *) v3,
699 (real *) w1, (real *) w2, (real *) w3, *n);
700 CUDA_CHECK(cudaGetLastError());
701 }
702
703
704 /*
705 * Reduction buffers, owned by the device layer and released
706 * on device teardown (cuda_buffer_free_all in cuda_finalize)
707 */
708 cuda_buffer_t redbuf = CUDA_BUFFER_INIT_SYMM;
709 cuda_buffer_t redbuf_xp = CUDA_BUFFER_INIT_SYMM;
710
714 void cuda_redbuf_check_alloc(int nb) {
715 cuda_buffer_reserve(&redbuf, (nb + 1) * sizeof(real));
716 }
717
721 void cuda_redbuf_check_alloc_xp(int nb) {
722 cuda_buffer_reserve(&redbuf_xp, (nb + 1) * sizeof(real_xp));
723 }
724
728 void cuda_global_reduce_add(real * bufred, void * bufred_d,
729 int n, const cudaStream_t stream) {
730
731#ifdef HAVE_NCCL
732 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
733 DEVICE_NCCL_SUM, stream);
734 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real)*n,
735 cudaMemcpyDeviceToHost, stream));
736 cudaStreamSynchronize(stream);
737#elif HAVE_NVSHMEM
738 if (sizeof(real) == sizeof(float)) {
739 nvshmemx_float_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
740 (float *) bufred_d,
741 (float *) bufred_d, n, stream);
742 }
743 else if (sizeof(real) == sizeof(double)) {
744 nvshmemx_double_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
745 (double *) bufred_d,
746 (double *) bufred_d, n, stream);
747
748 }
749 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
750 sizeof(real)*n, cudaMemcpyDeviceToHost, stream));
751 cudaStreamSynchronize(stream);
752#elif HAVE_DEVICE_MPI
753 cudaStreamSynchronize(stream);
754 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_SUM);
755#else
756 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real),
757 cudaMemcpyDeviceToHost, stream));
758 cudaStreamSynchronize(stream);
759#endif
760 }
761
765 void cuda_global_reduce_add_xp(real_xp * bufred, void * bufred_d,
766 int n, const cudaStream_t stream) {
767
768#ifdef HAVE_NCCL
769 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real_xp),
770 DEVICE_NCCL_SUM, stream);
771 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real_xp)*n,
772 cudaMemcpyDeviceToHost, stream));
773 cudaStreamSynchronize(stream);
774#elif HAVE_NVSHMEM
775 if (sizeof(real_xp) == sizeof(float)) {
776 nvshmemx_float_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
777 (float *) bufred_d,
778 (float *) bufred_d, n, stream);
779 }
780 else if (sizeof(real_xp) == sizeof(double)) {
781 nvshmemx_double_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD,
782 (double *) bufred_d,
783 (double *) bufred_d, n, stream);
784
785 }
786 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
787 sizeof(real_xp)*n, cudaMemcpyDeviceToHost, stream));
788 cudaStreamSynchronize(stream);
789#elif HAVE_DEVICE_MPI
790 cudaStreamSynchronize(stream);
791 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real_xp), DEVICE_MPI_SUM);
792#else
793 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real_xp),
794 cudaMemcpyDeviceToHost, stream));
795 cudaStreamSynchronize(stream);
796#endif
797 }
798
802 void cuda_global_reduce_max(real * bufred, void * bufred_d,
803 int n, const cudaStream_t stream) {
804
805#ifdef HAVE_NCCL
806 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
807 DEVICE_NCCL_MAX, stream);
808 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real)*n,
809 cudaMemcpyDeviceToHost, stream));
810 cudaStreamSynchronize(stream);
811#elif HAVE_NVSHMEM
812 if (sizeof(real) == sizeof(float)) {
813 nvshmemx_float_max_reduce_on_stream(NVSHMEM_TEAM_WORLD,
814 (float *) bufred_d,
815 (float *) bufred_d, n, stream);
816 }
817 else if (sizeof(real) == sizeof(double)) {
818 nvshmemx_double_max_reduce_on_stream(NVSHMEM_TEAM_WORLD,
819 (double *) bufred_d,
820 (double *) bufred_d, n, stream);
821
822 }
823 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
824 sizeof(real)*n, cudaMemcpyDeviceToHost, stream));
825 cudaStreamSynchronize(stream);
826#elif HAVE_DEVICE_MPI
827 cudaStreamSynchronize(stream);
828 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MAX);
829#else
830 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real),
831 cudaMemcpyDeviceToHost, stream));
832 cudaStreamSynchronize(stream);
833#endif
834 }
835
839 void cuda_global_reduce_min(real * bufred, void * bufred_d,
840 int n, const cudaStream_t stream) {
841
842#ifdef HAVE_NCCL
843 device_nccl_allreduce(bufred_d, bufred_d, n, sizeof(real),
844 DEVICE_NCCL_MIN, stream);
845 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, sizeof(real)*n,
846 cudaMemcpyDeviceToHost, stream));
847 cudaStreamSynchronize(stream);
848#elif HAVE_NVSHMEM
849 if (sizeof(real) == sizeof(float)) {
850 nvshmemx_float_min_reduce_on_stream(NVSHMEM_TEAM_WORLD,
851 (float *) bufred_d,
852 (float *) bufred_d, n, stream);
853 }
854 else if (sizeof(real) == sizeof(double)) {
855 nvshmemx_double_min_reduce_on_stream(NVSHMEM_TEAM_WORLD,
856 (double *) bufred_d,
857 (double *) bufred_d, n, stream);
858
859 }
860 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d,
861 sizeof(real)*n, cudaMemcpyDeviceToHost, stream));
862 cudaStreamSynchronize(stream);
863#elif HAVE_DEVICE_MPI
864 cudaStreamSynchronize(stream);
865 device_mpi_allreduce(bufred_d, bufred, n, sizeof(real), DEVICE_MPI_MIN);
866#else
867 CUDA_CHECK(cudaMemcpyAsync(bufred, bufred_d, n*sizeof(real),
868 cudaMemcpyDeviceToHost, stream));
869 cudaStreamSynchronize(stream);
870#endif
871 }
872
873
878 real cuda_vlsc3(void *u, void *v, void *w, int *n, cudaStream_t stream) {
879
880 const dim3 nthrds(1024, 1, 1);
881 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
882 const int nb = ((*n) + 1024 - 1)/ 1024;
883
884 cuda_redbuf_check_alloc(nb);
885
886 vlsc3_kernel<real><<<nblcks, nthrds, 0, stream>>>
887 ((real *) u, (real *) v, (real *) w, (real *) redbuf.dev, *n);
888 CUDA_CHECK(cudaGetLastError());
889 reduce_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev, nb);
890 CUDA_CHECK(cudaGetLastError());
891
892 CUDA_CHECK(cudaMemcpyAsync(((real *) redbuf.host), redbuf.dev, sizeof(real),
893 cudaMemcpyDeviceToHost, stream));
894 cudaStreamSynchronize(stream);
895
896 return ((real *) redbuf.host)[0];
897 }
898
899
900
901
906 real_xp cuda_glsc3(void *a, void *b, void *c, int *n, cudaStream_t stream) {
907
908 const dim3 nthrds(1024, 1, 1);
909 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
910 const int nb = ((*n) + 1024 - 1)/ 1024;
911
912 cuda_redbuf_check_alloc_xp(nb);
913
914 if ( *n > 0) {
915 glsc3_kernel<real, real_xp><<<nblcks, nthrds, 0, stream>>>
916 ((real *) a, (real *) b, (real *) c, (real_xp *) redbuf_xp.dev, *n);
917 CUDA_CHECK(cudaGetLastError());
918 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
919 CUDA_CHECK(cudaGetLastError());
920 }
921 else {
922 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
923 }
924 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
925
926 return ((real_xp *) redbuf_xp.host)[0];
927 }
928
933 void cuda_glsc3_many(real_xp *h, void * w, void *v,void *mult, int *j, int *n,
934 cudaStream_t stream){
935 int pow2 = 1;
936 while(pow2 < (*j)){
937 pow2 = 2*pow2;
938 }
939 const int nt = 1024/pow2;
940 const dim3 nthrds(pow2, nt, 1);
941 const dim3 nblcks(((*n)+nt - 1)/nt, 1, 1);
942 const int nb = ((*n) + nt - 1)/nt;
943
944 cuda_redbuf_check_alloc_xp((*j)*nb);
945
946 if ( *n > 0) {
947 glsc3_many_kernel<real, real_xp><<<nblcks, nthrds, 0, stream>>>
948 ((const real *) w, (const real **) v,
949 (const real *)mult, (real_xp *)redbuf_xp.dev, *j, *n);
950 CUDA_CHECK(cudaGetLastError());
951 glsc3_reduce_kernel<real_xp>
952 <<<(*j), 1024, 0, stream>>>((real_xp *) redbuf_xp.dev, nb, *j);
953 CUDA_CHECK(cudaGetLastError());
954 }
955 else {
956 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
957 }
958 cuda_global_reduce_add_xp(h, redbuf_xp.dev, (*j), stream);
959 }
960
965 real_xp cuda_glsc2(void *a, void *b, int *n, cudaStream_t stream) {
966
967 const dim3 nthrds(1024, 1, 1);
968 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
969 const int nb = ((*n) + 1024 - 1)/ 1024;
970
971 cuda_redbuf_check_alloc_xp(nb);
972
973 if ( *n > 0) {
974 glsc2_kernel<real, real_xp>
975 <<<nblcks, nthrds, 0, stream>>>((real *) a,
976 (real *) b,
977 (real_xp *) redbuf_xp.dev, *n);
978 CUDA_CHECK(cudaGetLastError());
979 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
980 CUDA_CHECK(cudaGetLastError());
981 }
982 else {
983 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
984 }
985 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
986
987 return ((real_xp *) redbuf_xp.host)[0];
988 }
989
994 real_xp cuda_glsubnorm2(void *a, void *b, int *n, cudaStream_t stream) {
995
996 const dim3 nthrds(1024, 1, 1);
997 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
998 const int nb = ((*n) + 1024 - 1)/ 1024;
999
1000 cuda_redbuf_check_alloc_xp(nb);
1001
1002 if ( *n > 0) {
1003 glsubnorm2_kernel<real, real_xp>
1004 <<<nblcks, nthrds, 0, stream>>>((real *) a,
1005 (real *) b,
1006 (real_xp *) redbuf_xp.dev, *n);
1007 CUDA_CHECK(cudaGetLastError());
1008 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
1009 CUDA_CHECK(cudaGetLastError());
1010 }
1011 else {
1012 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1013 }
1014 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
1015
1016 return ((real_xp *) redbuf_xp.host)[0];
1017 }
1018
1023 real_xp cuda_glsum(void *a, int *n, cudaStream_t stream) {
1024 const dim3 nthrds(1024, 1, 1);
1025 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1026 const int nb = ((*n) + 1024 - 1)/ 1024;
1027
1028 cuda_redbuf_check_alloc_xp(nb);
1029 if ( *n > 0) {
1030 glsum_kernel<real, real_xp>
1031 <<<nblcks, nthrds, 0, stream>>>((real *) a,
1032 (real_xp *) redbuf_xp.dev, *n);
1033 CUDA_CHECK(cudaGetLastError());
1034 reduce_kernel<real_xp><<<1, 1024, 0, stream>>> ((real_xp *) redbuf_xp.dev, nb);
1035 CUDA_CHECK(cudaGetLastError());
1036 }
1037 else {
1038 CUDA_CHECK(cudaMemsetAsync(redbuf_xp.dev, 0, redbuf_xp.size, stream));
1039 }
1040
1041 cuda_global_reduce_add_xp(((real_xp *) redbuf_xp.host), redbuf_xp.dev, 1, stream);
1042
1043 return ((real_xp *) redbuf_xp.host)[0];
1044 }
1045
1050 real cuda_glmax(void *a, real *ninf, int *n, cudaStream_t stream) {
1051 const dim3 nthrds(1024, 1, 1);
1052 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1053 const int nb = ((*n) + 1024 - 1)/ 1024;
1054
1055 cuda_redbuf_check_alloc(nb);
1056 if ( *n > 0) {
1057 glmax_kernel<real>
1058 <<<nblcks, nthrds, 0, stream>>>((real *) a, *ninf,
1059 (real *) redbuf.dev, *n);
1060 CUDA_CHECK(cudaGetLastError());
1061 reduce_max_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev,
1062 *ninf, nb);
1063 CUDA_CHECK(cudaGetLastError());
1064 }
1065 else {
1066 int nel = (int) (redbuf.size / sizeof(real));
1067 cuda_rzero(redbuf.dev, &nel, stream);
1068 }
1069
1070 cuda_global_reduce_max(((real *) redbuf.host), redbuf.dev, 1, stream);
1071
1072 return ((real *) redbuf.host)[0];
1073 }
1074
1079 real cuda_glmin(void *a, real *pinf, int *n, cudaStream_t stream) {
1080 const dim3 nthrds(1024, 1, 1);
1081 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1082 const int nb = ((*n) + 1024 - 1)/ 1024;
1083
1084 cuda_redbuf_check_alloc(nb);
1085 if ( *n > 0) {
1086 glmin_kernel<real>
1087 <<<nblcks, nthrds, 0, stream>>>((real *) a, *pinf,
1088 (real *) redbuf.dev, *n);
1089 CUDA_CHECK(cudaGetLastError());
1090 reduce_min_kernel<real><<<1, 1024, 0, stream>>> ((real *) redbuf.dev,
1091 *pinf, nb);
1092 CUDA_CHECK(cudaGetLastError());
1093 }
1094 else {
1095 int nel = (int) (redbuf.size / sizeof(real));
1096 cuda_rzero(redbuf.dev, &nel, stream);
1097 }
1098
1099 cuda_global_reduce_min(((real *) redbuf.host), redbuf.dev, 1, stream);
1100
1101 return ((real *) redbuf.host)[0];
1102 }
1103
1108 void cuda_absval(void *a, int *n, cudaStream_t stream) {
1109
1110 const dim3 nthrds(1024, 1, 1);
1111 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1112
1113 absval_kernel<real>
1114 <<<nblcks, nthrds,0, stream>>>((real *) a, * n);
1115 CUDA_CHECK(cudaGetLastError());
1116
1117 }
1118
1119 // ======================================================================== //
1120 // Point-wise operations.
1121
1126 void cuda_pwmax_vec2(void *a, void *b, int *n, cudaStream_t stream) {
1127
1128 const dim3 nthrds(1024, 1, 1);
1129 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1130
1131 pwmax_vec2_kernel<real><<<nblcks, nthrds, 0, stream>>>
1132 ((real *)a, (real *)b, *n);
1133 CUDA_CHECK(cudaGetLastError());
1134 }
1135
1140 void cuda_pwmax_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream) {
1141
1142 const dim3 nthrds(1024, 1, 1);
1143 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1144
1145 pwmax_vec3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1146 ((real *)a, (real *)b, (real *)c, *n);
1147 CUDA_CHECK(cudaGetLastError());
1148 }
1149
1154 void cuda_pwmax_sca2(void *a, real *c, int *n, cudaStream_t stream) {
1155
1156 const dim3 nthrds(1024, 1, 1);
1157 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1158
1159 pwmax_sca2_kernel<real><<<nblcks, nthrds, 0, stream>>>
1160 ((real *)a, *c, *n);
1161 CUDA_CHECK(cudaGetLastError());
1162 }
1163
1168 void cuda_pwmax_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream) {
1169
1170 const dim3 nthrds(1024, 1, 1);
1171 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1172
1173 pwmax_sca3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1174 ((real *)a, (real *)b, *c, *n);
1175 CUDA_CHECK(cudaGetLastError());
1176 }
1177
1182 void cuda_pwmin_vec2(void *a, void *b, int *n, cudaStream_t stream) {
1183
1184 const dim3 nthrds(1024, 1, 1);
1185 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1186
1187 pwmin_vec2_kernel<real><<<nblcks, nthrds, 0, stream>>>
1188 ((real *)a, (real *)b, *n);
1189 CUDA_CHECK(cudaGetLastError());
1190 }
1191
1196 void cuda_pwmin_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream) {
1197
1198 const dim3 nthrds(1024, 1, 1);
1199 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1200
1201 pwmin_vec3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1202 ((real *)a, (real *)b, (real *)c, *n);
1203 CUDA_CHECK(cudaGetLastError());
1204 }
1205
1210 void cuda_pwmin_sca2(void *a, real *c, int *n, cudaStream_t stream) {
1211
1212 const dim3 nthrds(1024, 1, 1);
1213 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1214
1215 pwmin_sca2_kernel<real><<<nblcks, nthrds, 0, stream>>>((real *)a, *c, *n);
1216 CUDA_CHECK(cudaGetLastError());
1217 }
1218
1223 void cuda_pwmin_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream) {
1224
1225 const dim3 nthrds(1024, 1, 1);
1226 const dim3 nblcks(((*n) + 1024 - 1) / 1024, 1, 1);
1227
1228 pwmin_sca3_kernel<real><<<nblcks, nthrds, 0, stream>>>
1229 ((real *)a, (real *)b, *c, *n);
1230 CUDA_CHECK(cudaGetLastError());
1231 }
1232
1233 // ======================================================================== //
1234
1238 void cuda_iadd(void *a, int *c, int *n, cudaStream_t stream) {
1239
1240 const dim3 nthrds(1024, 1, 1);
1241 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
1242
1243 cadd_kernel<int><<<nblcks, nthrds, 0, stream>>>
1244 ((int *) a, *c, *n);
1245 CUDA_CHECK(cudaGetLastError());
1246
1247 }
1248
1249} /* 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_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
CUDA_BUFFER_INIT_SYMM
#define CUDA_BUFFER_INIT_SYMM
Definition buffer.h:63
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
CUDA_CHECK
#define CUDA_CHECK(err)
Definition check.h:6
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:839
cuda_pwmax_sca3
void cuda_pwmax_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream)
Definition math.cu:1168
cuda_sqrt_inplace
void cuda_sqrt_inplace(void *a, int *n, cudaStream_t strm)
Definition math.cu:310
cuda_addcol4
void cuda_addcol4(void *a, void *b, void *c, void *d, int *n, cudaStream_t strm)
Definition math.cu:641
cuda_add2s2
void cuda_add2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm)
Definition math.cu:415
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:671
cuda_add3s2
void cuda_add3s2(void *a, void *b, void *c, real *c1, real *c2, int *n, cudaStream_t strm)
Definition math.cu:465
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:728
cuda_addcol3
void cuda_addcol3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:627
cuda_pwmin_vec2
void cuda_pwmin_vec2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:1182
cuda_pwmax_vec3
void cuda_pwmax_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream)
Definition math.cu:1140
cuda_power
void cuda_power(void *ap, void *a, real *p, int *n, cudaStream_t strm)
Definition math.cu:324
cuda_glmin
real cuda_glmin(void *a, real *pinf, int *n, cudaStream_t stream)
Definition math.cu:1079
cuda_cmult2
void cuda_cmult2(void *a, void *b, real *c, int *n, cudaStream_t strm)
Definition math.cu:225
cuda_add2s1
void cuda_add2s1(void *a, void *b, real *c1, int *n, cudaStream_t strm)
Definition math.cu:399
cuda_col3
void cuda_col3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:570
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:432
cuda_sub3
void cuda_sub3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:613
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:191
cuda_glsc3
real_xp cuda_glsc3(void *a, void *b, void *c, int *n, cudaStream_t stream)
Definition math.cu:906
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:529
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:765
cuda_pwmax_sca2
void cuda_pwmax_sca2(void *a, real *c, int *n, cudaStream_t stream)
Definition math.cu:1154
cuda_col2
void cuda_col2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:557
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:162
cuda_add2
void cuda_add2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:354
cuda_redbuf_check_alloc
void cuda_redbuf_check_alloc(int nb)
Definition math.cu:714
cuda_cmult
void cuda_cmult(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:212
cuda_add4
void cuda_add4(void *a, void *b, void *c, void *d, int *n, cudaStream_t strm)
Definition math.cu:383
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:176
cuda_sub2
void cuda_sub2(void *a, void *b, int *n, cudaStream_t strm)
Definition math.cu:599
cuda_redbuf_check_alloc_xp
void cuda_redbuf_check_alloc_xp(int nb)
Definition math.cu:721
cuda_invcol3
void cuda_invcol3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:543
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:688
cuda_glsc2
real_xp cuda_glsc2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:965
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:878
cuda_pwmin_sca3
void cuda_pwmin_sca3(void *a, void *b, real *c, int *n, cudaStream_t stream)
Definition math.cu:1223
cuda_cadd2
void cuda_cadd2(void *a, void *b, real *c, int *n, cudaStream_t strm)
Definition math.cu:280
cuda_addsqr2s2
void cuda_addsqr2s2(void *a, void *b, real *c1, int *n, cudaStream_t strm)
Definition math.cu:449
redbuf_xp
cuda_buffer_t redbuf_xp
Definition math.cu:709
cuda_glsum
real_xp cuda_glsum(void *a, int *n, cudaStream_t stream)
Definition math.cu:1023
cuda_add3
void cuda_add3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:369
cuda_addcol3s2
void cuda_addcol3s2(void *a, void *b, void *c, real *s, int *n, cudaStream_t strm)
Definition math.cu:656
cuda_cwrap
void cuda_cwrap(void *a, real *min_val, real *max_val, int *n, cudaStream_t strm)
Definition math.cu:295
cuda_pwmin_sca2
void cuda_pwmin_sca2(void *a, real *c, int *n, cudaStream_t stream)
Definition math.cu:1210
cuda_cfill
void cuda_cfill(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:338
redbuf
cuda_buffer_t redbuf
Definition math.cu:708
cuda_absval
void cuda_absval(void *a, int *n, cudaStream_t stream)
Definition math.cu:1108
cuda_rzero
void cuda_rzero(void *a, int *n, cudaStream_t strm)
Definition math.cu:205
cuda_cdiv2
void cuda_cdiv2(void *a, void *b, real *c, int *n, cudaStream_t strm)
Definition math.cu:252
cuda_cdiv
void cuda_cdiv(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:239
cuda_global_reduce_max
void cuda_global_reduce_max(real *bufred, void *bufred_d, int n, const cudaStream_t stream)
Definition math.cu:802
cuda_iadd
void cuda_iadd(void *a, int *c, int *n, cudaStream_t stream)
Definition math.cu:1238
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:1126
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:933
cuda_glsubnorm2
real_xp cuda_glsubnorm2(void *a, void *b, int *n, cudaStream_t stream)
Definition math.cu:994
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:499
cuda_invcol1
void cuda_invcol1(void *a, int *n, cudaStream_t strm)
Definition math.cu:516
cuda_radd
void cuda_radd(void *a, real *c, int *n, cudaStream_t strm)
Definition math.cu:266
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:482
cuda_glmax
real cuda_glmax(void *a, real *ninf, int *n, cudaStream_t stream)
Definition math.cu:1050
cuda_subcol3
void cuda_subcol3(void *a, void *b, void *c, int *n, cudaStream_t strm)
Definition math.cu:584
cuda_pwmin_vec3
void cuda_pwmin_vec3(void *a, void *b, void *c, int *n, cudaStream_t stream)
Definition math.cu:1196
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