Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
ax_helm.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 <string.h>
36#include <stdlib.h>
37#include <stdio.h>
38#include "ax_helm_kernel.h"
39#include "elem_block_tune.h"
40#include <device/device_config.h>
41#include <device/cuda/check.h>
42
43extern "C" {
44 #include <common/neko_log.h>
45}
46
47template < const int>
48int tune(void *w, void *u, void *dx, void *dy, void *dz,
49 void *dxt, void *dyt, void *dzt, void *h1,
50 void *g11, void *g22, void *g33, void *g12,
51 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel);
52
53template < const int>
54int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
55 void *dxt, void *dyt, void *dzt, void *h1,
56 void *g11, void *g22, void *g33, void *g12,
57 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel);
58
59extern "C" {
60
64 void cuda_ax_helm(void *w, void *u, void *dx, void *dy, void *dz,
65 void *dxt, void *dyt, void *dzt, void *h1,
66 void *g11, void *g22, void *g33, void *g12,
67 void *g13, void *g23, int *nelv, int *lx) {
68
69 static int autotune[17] = { 0 };
70 /* Elements per block candidate chosen by the autotuner, see
71 elem_block<> in elem_block.h */
72 static int autotune_eb[17] = { 0 };
73 /* chunk candidate chosen for the 1d variant */
74 static int autotune_ch[17] = { 0 };
75
76 const dim3 nthrds_1d(1024, 1, 1);
77 const dim3 nblcks_1d((*nelv), 1, 1);
78 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
79
80#define CASE_1D(LX, C) \
81 ax_helm_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
82 <<<nblcks_1d, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
83 ((real *) w, (real *) u, \
84 (real *) dx, (real *) dy, (real *) dz, \
85 (real *) dxt, (real *) dyt, (real *) dzt, (real *) h1,\
86 (real *) g11, (real *) g22, (real *) g33, \
87 (real *) g12, (real *) g13, (real *) g23); \
88 CUDA_CHECK(cudaGetLastError());
89
90/* Runtime dispatch onto the tuned chunk candidate */
91#define CASE_1D_SEL(LX, SEL) \
92 switch (SEL) { \
93 case 0: CASE_1D(LX, 0); break; \
94 case 1: CASE_1D(LX, 1); break; \
95 case 2: CASE_1D(LX, 2); break; \
96 default: CASE_1D(LX, 3); break; \
97 }
98
99#define CASE_KSTEP(LX, C) \
100 ax_helm_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
101 <<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
102 ((real *) w, (real *) u, \
103 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
104 (real *) g11, (real *) g22, (real *) g33, \
105 (real *) g12, (real *) g13, (real *) g23, *nelv); \
106 CUDA_CHECK(cudaGetLastError());
107
108#define CASE_KSTEP_PADDED(LX, C) \
109 ax_helm_kernel_kstep_padded<real, LX, NEKO_EB(LX, C)> \
110 <<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
111 ((real *) w, (real *) u, \
112 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
113 (real *) g11, (real *) g22, (real *) g33, \
114 (real *) g12, (real *) g13, (real *) g23, *nelv); \
115 CUDA_CHECK(cudaGetLastError());
116
117/* Runtime dispatch onto the tuned elements per block candidate */
118#define CASE_KSTEP_SEL(LX, SEL) \
119 switch (SEL) { \
120 case 0: CASE_KSTEP(LX, 0); break; \
121 case 1: CASE_KSTEP(LX, 1); break; \
122 default: CASE_KSTEP(LX, 2); break; \
123 }
124
125#define CASE_KSTEP_PADDED_SEL(LX, SEL) \
126 switch (SEL) { \
127 case 0: CASE_KSTEP_PADDED(LX, 0); break; \
128 case 1: CASE_KSTEP_PADDED(LX, 1); break; \
129 default: CASE_KSTEP_PADDED(LX, 2); break; \
130 }
131
132#define CASE(LX) \
133 case LX: \
134 if(autotune[LX] == 0 ) { \
135 autotune[LX]=tune<LX>( w, u, \
136 dx, dy, dz, \
137 dxt, dyt, dzt,h1, \
138 g11, g22, g33, \
139 g12, g13, g23, nelv, lx, \
140 &autotune_eb[LX], &autotune_ch[LX]); \
141 } else if (autotune[LX] == 1 ) { \
142 CASE_1D_SEL(LX, autotune_ch[LX]); \
143 } else if (autotune[LX] == 2 ) { \
144 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
145 } \
146 break
147
148#define CASE_PADDED(LX) \
149 case LX: \
150 if(autotune[LX] == 0 ) { \
151 autotune[LX]=tune_padded<LX>(w, u, \
152 dx, dy, dz, \
153 dxt, dyt, dzt,h1, \
154 g11, g22, g33, \
155 g12, g13, g23,nelv,lx, \
156 &autotune_eb[LX], &autotune_ch[LX]); \
157 } else if (autotune[LX] == 1 ) { \
158 CASE_1D_SEL(LX, autotune_ch[LX]); \
159 } else if (autotune[LX] == 2 ) { \
160 CASE_KSTEP_PADDED_SEL(LX, autotune_eb[LX]); \
161 } \
162 break
163
164/*
165 * High order cases have no 1d variant to compare against (its shared memory
166 * footprint grows as LX^3), so they are not tuned and keep one element per
167 * block, i.e. candidate 0
168 */
169#define CASE_LARGE(LX) \
170 case LX: \
171 CASE_KSTEP(LX, 0); \
172 break
173
174#define CASE_LARGE_PADDED(LX) \
175 case LX: \
176 CASE_KSTEP_PADDED(LX, 0); \
177 break
178
179 if ((*lx) < 12) {
180 switch(*lx) {
181 CASE(2);
182 CASE(3);
183 CASE_PADDED(4);
184 CASE(5);
185 CASE(6);
186 CASE(7);
187 CASE_PADDED(8);
188 CASE(9);
189 CASE(10);
190 CASE(11);
191 default:
192 {
193 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
194 exit(1);
195 }
196 }
197 }
198 else {
199 switch(*lx) {
200 CASE_LARGE(12);
201 CASE_LARGE(13);
202 CASE_LARGE(14);
203 CASE_LARGE(15);
204 CASE_LARGE_PADDED(16);
205 default:
206 {
207 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
208 exit(1);
209 }
210 }
211 }
212 }
213
217 void cuda_ax_helm_vector(void *au, void *av, void *aw,
218 void *u, void *v, void *w,
219 void *dx, void *dy, void *dz,
220 void *dxt, void *dyt, void *dzt,
221 void *h1, void *g11, void *g22,
222 void *g33, void *g12, void *g13,
223 void *g23, int *nelv, int *lx) {
224
225 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
226
227/*
228 * The vector kernels are not swept by the autotuner, see the note on
229 * NEKO_AX_HELM_VECTOR_EB_C in ax_helm_kernel.h
230 */
231#define AX_VEC_EB(LX) (elem_block<LX, NEKO_AX_HELM_VECTOR_EB_C>::value)
232#define AX_VEC_NTHRDS(LX) dim3((LX), (LX), AX_VEC_EB(LX))
233#define AX_VEC_NBLCKS(LX) \
234 dim3(((*nelv) + AX_VEC_EB(LX) - 1)/AX_VEC_EB(LX), 1, 1)
235
236#define CASE_VECTOR_KSTEP(LX) \
237 ax_helm_kernel_vector_kstep<real, LX, AX_VEC_EB(LX)> \
238 <<<AX_VEC_NBLCKS(LX), AX_VEC_NTHRDS(LX), 0, stream>>> \
239 ((real *) au, (real *) av, (real *) aw, \
240 (real *) u, (real *) v, (real *) w, \
241 (real *) dx, (real *) dy, (real *) dz, \
242 (real *) h1, (real *) g11, (real *) g22, \
243 (real *) g33, (real *) g12, (real *) g13, \
244 (real *) g23, *nelv); \
245 CUDA_CHECK(cudaGetLastError());
246
247#define CASE_VECTOR_KSTEP_PADDED(LX) \
248 ax_helm_kernel_vector_kstep_padded<real, LX, AX_VEC_EB(LX)> \
249 <<<AX_VEC_NBLCKS(LX), AX_VEC_NTHRDS(LX), 0, stream>>> \
250 ((real *) au, (real *) av, (real *) aw, \
251 (real *) u, (real *) v, (real *) w, \
252 (real *) dx, (real *) dy, (real *) dz, \
253 (real *) h1, (real *) g11, (real *) g22, \
254 (real *) g33, (real *) g12, (real *) g13, \
255 (real *) g23, *nelv); \
256 CUDA_CHECK(cudaGetLastError());
257
258#define CASE_VECTOR(LX) \
259 case LX: \
260 CASE_VECTOR_KSTEP(LX); \
261 break
262
263#define CASE_VECTOR_PADDED(LX) \
264 case LX: \
265 CASE_VECTOR_KSTEP_PADDED(LX); \
266 break
267
268 switch(*lx) {
269 CASE_VECTOR(2);
270 CASE_VECTOR(3);
271 CASE_VECTOR_PADDED(4);
272 CASE_VECTOR(5);
273 CASE_VECTOR(6);
274 CASE_VECTOR(7);
275 CASE_VECTOR_PADDED(8);
276 CASE_VECTOR(9);
277 CASE_VECTOR(10);
278 CASE_VECTOR(11);
279 CASE_VECTOR(12);
280 CASE_VECTOR(13);
281 CASE_VECTOR(14);
282 CASE_VECTOR(15);
283 CASE_VECTOR_PADDED(16);
284 default:
285 {
286 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
287 exit(1);
288 }
289 }
290 }
291
295 void cuda_ax_helm_vector_part2(void *au, void *av, void *aw,
296 void *u, void *v, void *w,
297 void *h2, void *B, int *n) {
298
299 const dim3 nthrds(1024, 1, 1);
300 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
301 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
302
303 ax_helm_kernel_vector_part2<real>
304 <<<nblcks, nthrds, 0, stream>>> ((real *) au, (real *) av, (real *) aw,
305 (real *) u, (real *) v, (real *) w,
306 (real *) h2, (real *) B, *n);
307 }
308}
309
310template < const int LX >
311int tune(void *w, void *u, void *dx, void *dy, void *dz,
312 void *dxt, void *dyt, void *dzt, void *h1,
313 void *g11, void *g22, void *g33, void *g12,
314 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel) {
315 cudaEvent_t start,stop;
316 float time1[NEKO_CHUNKS_CANDIDATES];
317 int best1 = 0;
318 float time2[NEKO_EB_CANDIDATES];
319 const int rounds = neko_tune_rounds();
320 const int iters = neko_tune_iters();
321 const int sweep = neko_eb_sweep();
322 int best = 0;
323 int retval;
324
325 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
326 time2[c] = NEKO_TUNE_INIT;
327 }
328 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
329 time1[c] = NEKO_TUNE_INIT;
330 }
331
332 const dim3 nthrds_1d(1024, 1, 1);
333 const dim3 nblcks_1d((*nelv), 1, 1);
334 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
335
336 char *env_value = NULL;
337 char neko_log_buf[80];
338
339 env_value=getenv("NEKO_AUTOTUNE");
340
341 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
342 log_section(neko_log_buf);
343
344 *eb_sel = 0;
345 *ch_sel = 0;
346
347 if(env_value) {
348 if( !strcmp(env_value,"1D") ) {
349 *ch_sel = neko_chunks_env();
350 CASE_1D_SEL(LX, *ch_sel);
351 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
352 NEKO_CHUNKS_SEL(LX, *ch_sel));
353 log_message(neko_log_buf);
354 log_end_section();
355 return 1;
356 } else if( !strcmp(env_value,"KSTEP") ) {
357 const int c = neko_eb_env();
358 *eb_sel = c;
359 CASE_KSTEP_SEL(LX, c);
360 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
361 NEKO_EB_SEL(LX, c));
362 log_message(neko_log_buf);
363 log_end_section();
364 return 2;
365 } else {
366 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
367 log_error(neko_log_buf);
368 }
369 }
370
371 cudaEventCreate(&start);
372 cudaEventCreate(&stop);
373
374 /* Warm every variant before timing anything: each specialisation has to
375 be resident and the clocks at steady state, or whichever variant is
376 timed first is measured on a colder part */
377 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
378 CASE_1D(LX, 0);
379 CASE_1D(LX, 1);
380 CASE_1D(LX, 2);
381 CASE_1D(LX, 3);
382 CASE_KSTEP(LX, 0);
383 if (sweep) {
384 CASE_KSTEP(LX, 1);
385 CASE_KSTEP(LX, 2);
386 }
387 }
388
389 /* Interleaved rounds, best time per variant: timing them one after another
390 in a fixed order lets clock drift bias the comparison by position */
391 for (int r = 0; r < rounds; r++) {
392 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
393 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
394 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
395 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
396 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 0, iters);
397 if (sweep) {
398 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 1, iters);
399 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 2, iters);
400 }
401 }
402
403 NEKO_TUNE_LOG(LX, time1, time2);
404
405 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
406 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
407 *eb_sel = best;
408 *ch_sel = best1;
409
410 if (time1[best1] < time2[best]) {
411 retval = 1;
412 } else {
413 retval = 2;
414 }
415
416 cudaEventDestroy(start);
417 cudaEventDestroy(stop);
418
419 /* The tuner stands in for a real Ax evaluation, and the variants do not
420 sum in the same order, so leave the output of the chosen kernel in w */
421 if (retval == 1) {
422 CASE_1D_SEL(LX, best1);
423 } else {
424 CASE_KSTEP_SEL(LX, best);
425 }
426
427 if (retval == 1) {
428 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
429 NEKO_CHUNKS_SEL(LX, best1));
430 } else {
431 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
432 NEKO_EB_SEL(LX, best));
433 }
434 log_message(neko_log_buf);
435 log_end_section();
436 return retval;
437}
438
439template < const int LX >
440int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
441 void *dxt, void *dyt, void *dzt, void *h1,
442 void *g11, void *g22, void *g33, void *g12,
443 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel) {
444 cudaEvent_t start, stop;
445 float time1[NEKO_CHUNKS_CANDIDATES];
446 int best1 = 0;
447 float time2[NEKO_EB_CANDIDATES];
448 const int rounds = neko_tune_rounds();
449 const int iters = neko_tune_iters();
450 const int sweep = neko_eb_sweep();
451 int best = 0;
452 int retval;
453
454 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
455 time2[c] = NEKO_TUNE_INIT;
456 }
457 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
458 time1[c] = NEKO_TUNE_INIT;
459 }
460
461 const dim3 nthrds_1d(1024, 1, 1);
462 const dim3 nblcks_1d((*nelv), 1, 1);
463 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
464
465 char *env_value = NULL;
466 char neko_log_buf[80];
467
468 env_value=getenv("NEKO_AUTOTUNE");
469
470 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
471 log_section(neko_log_buf);
472
473 *eb_sel = 0;
474 *ch_sel = 0;
475
476 if(env_value) {
477 if( !strcmp(env_value,"1D") ) {
478 *ch_sel = neko_chunks_env();
479 CASE_1D_SEL(LX, *ch_sel);
480 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
481 NEKO_CHUNKS_SEL(LX, *ch_sel));
482 log_message(neko_log_buf);
483 log_end_section();
484 return 1;
485 } else if( !strcmp(env_value,"KSTEP") ) {
486 const int c = neko_eb_env();
487 *eb_sel = c;
488 CASE_KSTEP_PADDED_SEL(LX, c);
489 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
490 NEKO_EB_SEL(LX, c));
491 log_message(neko_log_buf);
492 log_end_section();
493 return 2;
494 } else {
495 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
496 log_error(neko_log_buf);
497 }
498 }
499
500 cudaEventCreate(&start);
501 cudaEventCreate(&stop);
502
503 /* Warm every variant before timing anything: each specialisation has to
504 be resident and the clocks at steady state, or whichever variant is
505 timed first is measured on a colder part */
506 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
507 CASE_1D(LX, 0);
508 CASE_1D(LX, 1);
509 CASE_1D(LX, 2);
510 CASE_1D(LX, 3);
511 CASE_KSTEP_PADDED(LX, 0);
512 if (sweep) {
513 CASE_KSTEP_PADDED(LX, 1);
514 CASE_KSTEP_PADDED(LX, 2);
515 }
516 }
517
518 /* Interleaved rounds, best time per variant: timing them one after another
519 in a fixed order lets clock drift bias the comparison by position */
520 for (int r = 0; r < rounds; r++) {
521 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
522 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
523 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
524 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
525 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 0, iters);
526 if (sweep) {
527 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 1, iters);
528 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 2, iters);
529 }
530 }
531
532 NEKO_TUNE_LOG(LX, time1, time2);
533
534 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
535 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
536 *eb_sel = best;
537 *ch_sel = best1;
538
539 if (time1[best1] < time2[best]) {
540 retval=1;
541 } else {
542 retval=2;
543 }
544
545 cudaEventDestroy(start);
546 cudaEventDestroy(stop);
547
548 /* Leave the output of the chosen kernel in w, see tune() */
549 if (retval == 1) {
550 CASE_1D_SEL(LX, best1);
551 } else {
552 CASE_KSTEP_PADDED_SEL(LX, best);
553 }
554
555 if (retval == 1) {
556 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
557 NEKO_CHUNKS_SEL(LX, best1));
558 } else {
559 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
560 NEKO_EB_SEL(LX, best));
561 }
562 log_message(neko_log_buf);
563 log_end_section();
564 return retval;
565}
CASE_VECTOR
#define CASE_VECTOR(LX)
CASE_KSTEP_SEL
#define CASE_KSTEP_SEL(LX, SEL)
cuda_ax_helm_vector_part2
void cuda_ax_helm_vector_part2(void *au, void *av, void *aw, void *u, void *v, void *w, void *h2, void *B, int *n)
Definition ax_helm.cu:295
CASE
#define CASE(LX)
CASE_1D_SEL
#define CASE_1D_SEL(LX, SEL)
tune
int tune(void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel)
Definition ax_helm.cu:311
cuda_ax_helm
void cuda_ax_helm(void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.cu:64
CASE_PADDED
#define CASE_PADDED(LX)
cuda_ax_helm_vector
void cuda_ax_helm_vector(void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.cu:217
CASE_KSTEP
#define CASE_KSTEP(LX, C)
CASE_KSTEP_PADDED_SEL
#define CASE_KSTEP_PADDED_SEL(LX, SEL)
CASE_LARGE
#define CASE_LARGE(LX)
CASE_LARGE_PADDED
#define CASE_LARGE_PADDED(LX)
tune_padded
int tune_padded(void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel)
Definition ax_helm.cu:440
CASE_1D
#define CASE_1D(LX, C)
CASE_VECTOR_PADDED
#define CASE_VECTOR_PADDED(LX)
CASE_KSTEP_PADDED
#define CASE_KSTEP_PADDED(LX, C)
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
aw
__global__ void T *__restrict__ T *__restrict__ aw
Definition ax_helm_full_kernel.h:45
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
i
const int i
Definition ax_helm_full_kernel.h:95
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_full_kernel.h:49
av
__global__ void T *__restrict__ av
Definition ax_helm_full_kernel.h:44
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition ax_helm_full_kernel.h:51
h1
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1
Definition ax_helm_full_kernel.h:52
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_full_kernel.h:50
nelv
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const int nelv
Definition ax_helm_kernel.h:177
g23
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g23
Definition ax_helm_kernel.h:176
g22
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g22
Definition ax_helm_kernel.h:172
g13
__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__ const T *__restrict__ const T *__restrict__ g13
Definition ax_helm_kernel.h:175
g12
__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__ const T *__restrict__ g12
Definition ax_helm_kernel.h:174
g33
__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__ g33
Definition ax_helm_kernel.h:173
g11
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g11
Definition ax_helm_kernel.h:171
dzt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
Definition cdtp_kernel.h:114
dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
Definition cdtp_kernel.h:113
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition cdtp_kernel.h:112
NEKO_CHUNKS_CANDIDATES
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
NEKO_EB_CANDIDATES
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
NEKO_EB_SEL
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
NEKO_CHUNKS_SEL
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
NEKO_TUNE_TIME
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
Definition elem_block_tune.h:119
neko_eb_env
static int neko_eb_env()
Definition elem_block_tune.h:79
neko_tune_rounds
static int neko_tune_rounds()
Definition elem_block_tune.h:102
NEKO_TUNE_LOG
#define NEKO_TUNE_LOG(LX, T1, T2)
Definition elem_block_tune.h:138
NEKO_TUNE_INIT
#define NEKO_TUNE_INIT
Definition elem_block_tune.h:65
NEKO_TUNE_BEST
#define NEKO_TUNE_BEST(T, BEST, N)
Definition elem_block_tune.h:130
neko_tune_iters
static int neko_tune_iters()
Definition elem_block_tune.h:110
neko_chunks_env
static int neko_chunks_env()
Definition elem_block_tune.h:91
neko_eb_sweep
static int neko_eb_sweep()
Definition elem_block_tune.h:68
NEKO_TUNE_WARMUP
#define NEKO_TUNE_WARMUP
Definition elem_block_tune.h:66
device_config.h
real
double real
Definition device_config.h:12
ax_helm_kernel.h
elem_block_tune.h
neko_log.h
log_error
void log_error(char *msg)
log_message
void log_message(char *msg)
log_end_section
void log_end_section()
log_section
void log_section(char *msg)