Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
ax_helm.hip
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 <hip/hip_runtime.h>
39#include <device/device_config.h>
40#include <device/hip/check.h>
41#include "ax_helm_kernel.h"
42#include "elem_block_tune.h"
43
44extern "C" {
45 #include <common/neko_log.h>
46}
47
48template < const int>
49int tune(void *w, void *u, void *dx, void *dy, void *dz,
50 void *dxt, void *dyt, void *dzt, void *h1,
51 void *g11, void *g22, void *g33, void *g12,
52 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel);
53
54template < const int>
55int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
56 void *dxt, void *dyt, void *dzt, void *h1,
57 void *g11, void *g22, void *g33, void *g12,
58 void *g13, void *g23, int *nelv, int *lx, int *eb_sel,
59 int *ch_sel);
60
61template < const int>
62int tune_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
63 void *dx, void *dy, void *dz, void *h1,
64 void *g11, void *g22, void *g33, void *g12,
65 void *g13, void *g23, int *nelv, int *lx);
66
67template < const int>
68int tune_vector_padded(void *au, void *av, void *aw,
69 void *u, void *v, void *w,
70 void *dx, void *dy, void *dz, void *h1,
71 void *g11, void *g22, void *g33, void *g12,
72 void *g13, void *g23, int *nelv, int *lx);
73
74extern "C" {
75
79 void hip_ax_helm(void *w, void *u, void *dx, void *dy, void *dz,
80 void *dxt, void *dyt, void *dzt, void *h1,
81 void *g11, void *g22, void *g33, void *g12,
82 void *g13, void *g23, int *nelv, int *lx) {
83
84 static int autotune[13] = { 0 };
85 /* Elements per block candidate chosen by the autotuner, see
86 elem_block<> in elem_block.h */
87 static int autotune_eb[13] = { 0 };
88 /* chunk candidate chosen for the 1d variant */
89 static int autotune_ch[13] = { 0 };
90
91 const dim3 nthrds_1d(1024, 1, 1);
92 const dim3 nblcks_1d((*nelv), 1, 1);
93
94#define CASE_1D(LX, C) \
95 hipLaunchKernelGGL(HIP_KERNEL_NAME( \
96 ax_helm_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> ), \
97 nblcks_1d, NEKO_CHUNKS_NTHRDS(LX, C), 0, \
98 (hipStream_t) glb_cmd_queue, \
99 (real *) w, (real *) u, \
100 (real *) dx, (real *) dy, (real *) dz, \
101 (real *) dxt, (real *) dyt, (real *) dzt, (real *) h1, \
102 (real *) g11, (real *) g22, (real *) g33, \
103 (real *) g12, (real *) g13, (real *) g23); \
104 HIP_CHECK(hipGetLastError());
105
106/* Runtime dispatch onto the tuned chunk candidate */
107#define CASE_1D_SEL(LX, SEL) \
108 switch (SEL) { \
109 case 0: CASE_1D(LX, 0); break; \
110 case 1: CASE_1D(LX, 1); break; \
111 case 2: CASE_1D(LX, 2); break; \
112 default: CASE_1D(LX, 3); break; \
113 }
114
115#define CASE_KSTEP(LX, C) \
116 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
117 ax_helm_kernel_kstep<real, LX, NEKO_EB(LX, C)> ), \
118 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
119 (hipStream_t) glb_cmd_queue, \
120 (real *) w, (real *) u, \
121 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
122 (real *) g11, (real *) g22, (real *) g33, \
123 (real *) g12, (real *) g13, (real *) g23, *nelv); \
124 HIP_CHECK(hipGetLastError());
125
126#define CASE_KSTEP_PADDED(LX, C) \
127 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
128 ax_helm_kernel_kstep_padded<real, LX, NEKO_EB(LX, C)> ),\
129 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
130 (hipStream_t) glb_cmd_queue, \
131 (real *) w, (real *) u, \
132 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
133 (real *) g11, (real *) g22, (real *) g33, \
134 (real *) g12, (real *) g13, (real *) g23, *nelv); \
135 HIP_CHECK(hipGetLastError());
136
137/* Runtime dispatch onto the tuned elements per block candidate */
138#define CASE_KSTEP_SEL(LX, SEL) \
139 switch (SEL) { \
140 case 0: CASE_KSTEP(LX, 0); break; \
141 case 1: CASE_KSTEP(LX, 1); break; \
142 default: CASE_KSTEP(LX, 2); break; \
143 }
144
145#define CASE_KSTEP_PADDED_SEL(LX, SEL) \
146 switch (SEL) { \
147 case 0: CASE_KSTEP_PADDED(LX, 0); break; \
148 case 1: CASE_KSTEP_PADDED(LX, 1); break; \
149 default: CASE_KSTEP_PADDED(LX, 2); break; \
150 }
151
152#define CASE(LX) \
153 case LX: \
154 if(autotune[LX] == 0 ) { \
155 autotune[LX]=tune<LX>( w, u, \
156 dx, dy, dz, \
157 dxt, dyt, dzt,h1, \
158 g11, g22, g33, \
159 g12, g13, g23, nelv, lx, \
160 &autotune_eb[LX], &autotune_ch[LX]); \
161 } else if (autotune[LX] == 1 ) { \
162 CASE_1D_SEL(LX, autotune_ch[LX]); \
163 } else if (autotune[LX] == 2 ) { \
164 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
165 } \
166 break
167
168#define CASE_PADDED(LX) \
169 case LX: \
170 if(autotune[LX] == 0 ) { \
171 autotune[LX]=tune_padded<LX>(w, u, \
172 dx, dy, dz, \
173 dxt, dyt, dzt,h1, \
174 g11, g22, g33, \
175 g12, g13, g23,nelv,lx, \
176 &autotune_eb[LX], &autotune_ch[LX]); \
177 } else if (autotune[LX] == 1 ) { \
178 CASE_1D_SEL(LX, autotune_ch[LX]); \
179 } else if (autotune[LX] == 2 ) { \
180 CASE_KSTEP_PADDED_SEL(LX, autotune_eb[LX]); \
181 } \
182 break
183
184/*
185 * High order cases have no 1d variant to compare against (its LDS footprint
186 * grows as LX^3), so they are not tuned and keep one element per block,
187 * i.e. candidate 0
188 */
189#define CASE_LARGE(LX) \
190 case LX: \
191 CASE_KSTEP(LX, 0); \
192 break
193
194#define CASE_LARGE_PADDED(LX) \
195 case LX: \
196 CASE_KSTEP_PADDED(LX, 0); \
197 break
198
199 if ((*lx) < 13) {
200 switch(*lx) {
201 CASE(2);
202 CASE(3);
203 CASE(4);
204 CASE(5);
205 CASE(6);
206 CASE(7);
207 CASE_PADDED(8);
208 CASE(9);
209 CASE(10);
210 CASE(11);
211 CASE(12);
212 default:
213 {
214 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
215 exit(1);
216 }
217 }
218 }
219 else {
220 switch(*lx) {
221 CASE_LARGE(13);
222 CASE_LARGE(14);
223 CASE_LARGE(15);
224 CASE_LARGE_PADDED(16);
225 default:
226 {
227 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
228 exit(1);
229 }
230 }
231 }
232 }
236 void hip_ax_helm_vector(void *au, void *av, void *aw,
237 void *u, void *v, void *w,
238 void *dx, void *dy, void *dz,
239 void *dxt, void *dyt, void *dzt,
240 void *h1, void *g11, void *g22,
241 void *g33, void *g12, void *g13,
242 void *g23, int *nelv, int *lx) {
243
244 /* 0 = untuned, otherwise the chosen candidate plus one. Unlike the
245 scalar Ax there is no 1d variant to fall back on, so the tuner only
246 sweeps the elements per block candidates */
247 static int autotune_v[17] = { 0 };
248
249#define CASE_VECTOR_KSTEP(LX, C) \
250 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
251 ax_helm_kernel_vector_kstep<real, LX, \
252 NEKO_EB(LX, C)> ), \
253 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
254 (hipStream_t) glb_cmd_queue, \
255 (real *) au, (real *) av, (real *) aw, \
256 (real *) u, (real *) v, (real *) w, \
257 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
258 (real *) g11, (real *) g22, (real *) g33, \
259 (real *) g12, (real *) g13, (real *) g23, *nelv); \
260 HIP_CHECK(hipGetLastError());
261
262#define CASE_VECTOR_KSTEP_PADDED(LX, C) \
263 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
264 ax_helm_kernel_vector_kstep_padded<real, LX, \
265 NEKO_EB(LX, C)> ), \
266 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
267 (hipStream_t) glb_cmd_queue, \
268 (real *) au, (real *) av, (real *) aw, \
269 (real *) u, (real *) v, (real *) w, \
270 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
271 (real *) g11, (real *) g22, (real *) g33, \
272 (real *) g12, (real *) g13, (real *) g23, *nelv); \
273 HIP_CHECK(hipGetLastError());
274
275#define CASE_VECTOR_SEL(LX, SEL) \
276 switch (SEL) { \
277 case 0: CASE_VECTOR_KSTEP(LX, 0); break; \
278 case 1: CASE_VECTOR_KSTEP(LX, 1); break; \
279 default: CASE_VECTOR_KSTEP(LX, 2); break; \
280 }
281
282#define CASE_VECTOR_PADDED_SEL(LX, SEL) \
283 switch (SEL) { \
284 case 0: CASE_VECTOR_KSTEP_PADDED(LX, 0); break; \
285 case 1: CASE_VECTOR_KSTEP_PADDED(LX, 1); break; \
286 default: CASE_VECTOR_KSTEP_PADDED(LX, 2); break; \
287 }
288
289#define CASE_VECTOR(LX) \
290 case LX: \
291 if (autotune_v[LX] == 0) { \
292 autotune_v[LX] = 1 + tune_vector<LX>(au, av, aw, u, v, w, \
293 dx, dy, dz, h1, \
294 g11, g22, g33, \
295 g12, g13, g23, nelv, lx); \
296 } else { \
297 CASE_VECTOR_SEL(LX, autotune_v[LX] - 1); \
298 } \
299 break
300
301#define CASE_VECTOR_PADDED(LX) \
302 case LX: \
303 if (autotune_v[LX] == 0) { \
304 autotune_v[LX] = 1 + tune_vector_padded<LX>(au, av, aw, u, v, w, \
305 dx, dy, dz, h1, \
306 g11, g22, g33, \
307 g12, g13, g23, nelv, lx); \
308 } else { \
309 CASE_VECTOR_PADDED_SEL(LX, autotune_v[LX] - 1); \
310 } \
311 break
312
313 switch(*lx) {
314 CASE_VECTOR(2);
315 CASE_VECTOR(3);
316 CASE_VECTOR_PADDED(4);
317 CASE_VECTOR(5);
318 CASE_VECTOR(6);
319 CASE_VECTOR(7);
320 CASE_VECTOR_PADDED(8);
321 CASE_VECTOR(9);
322 CASE_VECTOR(10);
323 CASE_VECTOR(11);
324 CASE_VECTOR(12);
325 CASE_VECTOR(13);
326 CASE_VECTOR(14);
327 CASE_VECTOR(15);
328 CASE_VECTOR_PADDED(16);
329 default:
330 {
331 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
332 exit(1);
333 }
334 }
335 }
336
340 void hip_ax_helm_vector_part2(void *au, void *av, void *aw,
341 void *u, void *v, void *w,
342 void *h2, void *B, int *n) {
343
344 const dim3 nthrds(1024, 1, 1);
345 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
346 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
347
348 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_vector_part2<real> ),
349 nblcks, nthrds, 0, stream,
350 (real *) au, (real *) av, (real *) aw,
351 (real *) u, (real *) v, (real *) w,
352 (real *) h2, (real *) B, *n);
353 }
354
355}
356
357template < const int LX >
358int tune(void *w, void *u, void *dx, void *dy, void *dz,
359 void *dxt, void *dyt, void *dzt, void *h1,
360 void *g11, void *g22, void *g33, void *g12,
361 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel) {
362 hipEvent_t start,stop;
363 float time1[NEKO_CHUNKS_CANDIDATES];
364 int best1 = 0;
365 float time2[NEKO_EB_CANDIDATES];
366 const int rounds = neko_tune_rounds();
367 const int iters = neko_tune_iters();
368 const int sweep = neko_eb_sweep();
369 int best = 0;
370 int retval;
371
372 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
373 time2[c] = NEKO_TUNE_INIT;
374 }
375 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
376 time1[c] = NEKO_TUNE_INIT;
377 }
378
379 const dim3 nthrds_1d(1024, 1, 1);
380 const dim3 nblcks_1d((*nelv), 1, 1);
381 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
382
383 char *env_value = NULL;
384 char neko_log_buf[80];
385
386 env_value=getenv("NEKO_AUTOTUNE");
387
388 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
389 log_section(neko_log_buf);
390
391 *eb_sel = 0;
392 *ch_sel = 0;
393
394 if(env_value) {
395 if( !strcmp(env_value,"1D") ) {
396 *ch_sel = neko_chunks_env();
397 CASE_1D_SEL(LX, *ch_sel);
398 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
399 NEKO_CHUNKS_SEL(LX, *ch_sel));
400 log_message(neko_log_buf);
401 log_end_section();
402 return 1;
403 } else if( !strcmp(env_value,"KSTEP") ) {
404 const int c = neko_eb_env();
405 *eb_sel = c;
406 CASE_KSTEP_SEL(LX, c);
407 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
408 NEKO_EB_SEL(LX, c));
409 log_message(neko_log_buf);
410 log_end_section();
411 return 2;
412 } else {
413 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
414 log_error(neko_log_buf);
415 }
416 }
417
418 HIP_CHECK(hipEventCreate(&start));
419 HIP_CHECK(hipEventCreate(&stop));
420
421 /* Warm every variant before timing anything: each specialisation has to
422 be resident and the clocks at steady state, or whichever variant is
423 timed first is measured on a colder part */
424 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
425 CASE_1D(LX, 0);
426 CASE_1D(LX, 1);
427 CASE_1D(LX, 2);
428 CASE_1D(LX, 3);
429 CASE_KSTEP(LX, 0);
430 if (sweep) {
431 CASE_KSTEP(LX, 1);
432 CASE_KSTEP(LX, 2);
433 }
434 }
435
436 /* Interleaved rounds, best time per variant: timing them one after another
437 in a fixed order lets clock drift bias the comparison by position */
438 for (int r = 0; r < rounds; r++) {
439 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
440 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
441 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
442 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
443 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 0, iters);
444 if (sweep) {
445 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 1, iters);
446 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 2, iters);
447 }
448 }
449
450 NEKO_TUNE_LOG(LX, time1, time2);
451
452 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
453 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
454 *eb_sel = best;
455 *ch_sel = best1;
456
457 if (time1[best1] < time2[best]) {
458 retval = 1;
459 } else {
460 retval = 2;
461 }
462
463 HIP_CHECK(hipEventDestroy(start));
464 HIP_CHECK(hipEventDestroy(stop));
465
466 /* The tuner stands in for a real Ax evaluation, and the variants do not
467 sum in the same order, so leave the output of the chosen kernel in w */
468 if (retval == 1) {
469 CASE_1D_SEL(LX, best1);
470 } else {
471 CASE_KSTEP_SEL(LX, best);
472 }
473
474 if (retval == 1) {
475 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
476 NEKO_CHUNKS_SEL(LX, best1));
477 } else {
478 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
479 NEKO_EB_SEL(LX, best));
480 }
481 log_message(neko_log_buf);
482 log_end_section();
483 return retval;
484}
485
486template < const int LX >
487int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
488 void *dxt, void *dyt, void *dzt, void *h1,
489 void *g11, void *g22, void *g33, void *g12,
490 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel) {
491 hipEvent_t start, stop;
492 float time1[NEKO_CHUNKS_CANDIDATES];
493 int best1 = 0;
494 float time2[NEKO_EB_CANDIDATES];
495 const int rounds = neko_tune_rounds();
496 const int iters = neko_tune_iters();
497 const int sweep = neko_eb_sweep();
498 int best = 0;
499 int retval;
500
501 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
502 time2[c] = NEKO_TUNE_INIT;
503 }
504 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
505 time1[c] = NEKO_TUNE_INIT;
506 }
507
508 const dim3 nthrds_1d(1024, 1, 1);
509 const dim3 nblcks_1d((*nelv), 1, 1);
510 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
511
512 char *env_value = NULL;
513 char neko_log_buf[80];
514
515 env_value=getenv("NEKO_AUTOTUNE");
516
517 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
518 log_section(neko_log_buf);
519
520 *eb_sel = 0;
521 *ch_sel = 0;
522
523 if(env_value) {
524 if( !strcmp(env_value,"1D") ) {
525 *ch_sel = neko_chunks_env();
526 CASE_1D_SEL(LX, *ch_sel);
527 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
528 NEKO_CHUNKS_SEL(LX, *ch_sel));
529 log_message(neko_log_buf);
530 log_end_section();
531 return 1;
532 } else if( !strcmp(env_value,"KSTEP") ) {
533 const int c = neko_eb_env();
534 *eb_sel = c;
535 CASE_KSTEP_PADDED_SEL(LX, c);
536 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
537 NEKO_EB_SEL(LX, c));
538 log_message(neko_log_buf);
539 log_end_section();
540 return 2;
541 } else {
542 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
543 log_error(neko_log_buf);
544 }
545 }
546
547 HIP_CHECK(hipEventCreate(&start));
548 HIP_CHECK(hipEventCreate(&stop));
549
550 /* Warm every variant before timing anything: each specialisation has to
551 be resident and the clocks at steady state, or whichever variant is
552 timed first is measured on a colder part */
553 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
554 CASE_1D(LX, 0);
555 CASE_1D(LX, 1);
556 CASE_1D(LX, 2);
557 CASE_1D(LX, 3);
558 CASE_KSTEP_PADDED(LX, 0);
559 if (sweep) {
560 CASE_KSTEP_PADDED(LX, 1);
561 CASE_KSTEP_PADDED(LX, 2);
562 }
563 }
564
565 /* Interleaved rounds, best time per variant: timing them one after another
566 in a fixed order lets clock drift bias the comparison by position */
567 for (int r = 0; r < rounds; r++) {
568 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
569 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
570 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
571 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
572 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 0, iters);
573 if (sweep) {
574 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 1, iters);
575 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 2, iters);
576 }
577 }
578
579 NEKO_TUNE_LOG(LX, time1, time2);
580
581 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
582 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
583 *eb_sel = best;
584 *ch_sel = best1;
585
586 if (time1[best1] < time2[best]) {
587 retval=1;
588 } else {
589 retval=2;
590 }
591
592 HIP_CHECK(hipEventDestroy(start));
593 HIP_CHECK(hipEventDestroy(stop));
594
595 /* Leave the output of the chosen kernel in w, see tune() */
596 if (retval == 1) {
597 CASE_1D_SEL(LX, best1);
598 } else {
599 CASE_KSTEP_PADDED_SEL(LX, best);
600 }
601
602 if (retval == 1) {
603 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
604 NEKO_CHUNKS_SEL(LX, best1));
605 } else {
606 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
607 NEKO_EB_SEL(LX, best));
608 }
609 log_message(neko_log_buf);
610 log_end_section();
611 return retval;
612}
613
614/*
615 * The vector Ax has no 1d variant, so these only sweep the elements per
616 * block candidates. Setting NEKO_AUTOTUNE to anything pins candidate 0,
617 * i.e. the one element per block shape these kernels have always had,
618 * which is the A/B baseline for the blocking.
619 */
620template < const int LX >
621int tune_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
622 void *dx, void *dy, void *dz, void *h1,
623 void *g11, void *g22, void *g33, void *g12,
624 void *g13, void *g23, int *nelv, int *lx) {
625 hipEvent_t start, stop;
626 float time[NEKO_EB_CANDIDATES];
627 const int rounds = neko_tune_rounds();
628 const int iters = neko_tune_iters();
629 const int sweep = neko_eb_sweep();
630 int best = 0;
631
632 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
633 time[c] = NEKO_TUNE_INIT;
634 }
635
636 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
637 char neko_log_buf[80];
638
639 sprintf(neko_log_buf, "Autotune Ax helm vector (lx: %d)", *lx);
640 log_section(neko_log_buf);
641
642 /*
643 * Default to one element per block, as on the CUDA backend. These kernels
644 * carry six T[LX] register arrays against the scalar kernels' two, so they
645 * reach the register file limit -- and start spilling to scratch -- well
646 * before a wider block can buy anything. Sweeping them by default was a
647 * regression. Set NEKO_EB_TUNE=1 to opt back into the sweep
648 */
649 if (getenv("NEKO_AUTOTUNE") ||
650 getenv("NEKO_EB") || !neko_eb_sweep()) {
651 const int c = neko_eb_env();
652 CASE_VECTOR_SEL(LX, c);
653 sprintf(neko_log_buf, "Vector Ax : %d elem/block (untuned)",
654 NEKO_EB_SEL(LX, c));
655 log_message(neko_log_buf);
656 log_end_section();
657 return c;
658 }
659
660 HIP_CHECK(hipEventCreate(&start));
661 HIP_CHECK(hipEventCreate(&stop));
662
663 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
664 CASE_VECTOR_KSTEP(LX, 0);
665 CASE_VECTOR_KSTEP(LX, 1);
666 CASE_VECTOR_KSTEP(LX, 2);
667 }
668
669 for (int r = 0; r < rounds; r++) {
670 NEKO_TUNE_TIME(time, CASE_VECTOR_KSTEP, LX, 0, iters);
671 NEKO_TUNE_TIME(time, CASE_VECTOR_KSTEP, LX, 1, iters);
672 NEKO_TUNE_TIME(time, CASE_VECTOR_KSTEP, LX, 2, iters);
673 }
674
675 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
676 sprintf(neko_log_buf, "eb=%-4d : %9.2f us/call",
677 NEKO_EB_SEL(LX, c), time[c] * 10.0);
678 log_message(neko_log_buf);
679 }
680
681 NEKO_TUNE_BEST(time, best, NEKO_EB_CANDIDATES);
682
683 HIP_CHECK(hipEventDestroy(start));
684 HIP_CHECK(hipEventDestroy(stop));
685
686 /* Leave the output of the chosen kernel in au/av/aw */
687 CASE_VECTOR_SEL(LX, best);
688
689 sprintf(neko_log_buf, "Chose : %d elem/block", NEKO_EB_SEL(LX, best));
690 log_message(neko_log_buf);
691 log_end_section();
692 return best;
693}
694
695template < const int LX >
696int tune_vector_padded(void *au, void *av, void *aw,
697 void *u, void *v, void *w,
698 void *dx, void *dy, void *dz, void *h1,
699 void *g11, void *g22, void *g33, void *g12,
700 void *g13, void *g23, int *nelv, int *lx) {
701 hipEvent_t start, stop;
702 float time[NEKO_EB_CANDIDATES];
703 const int rounds = neko_tune_rounds();
704 const int iters = neko_tune_iters();
705 const int sweep = neko_eb_sweep();
706 int best = 0;
707
708 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
709 time[c] = NEKO_TUNE_INIT;
710 }
711
712 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
713 char neko_log_buf[80];
714
715 sprintf(neko_log_buf, "Autotune Ax helm vector (lx: %d)", *lx);
716 log_section(neko_log_buf);
717
718 /*
719 * Default to one element per block, as on the CUDA backend. These kernels
720 * carry six T[LX] register arrays against the scalar kernels' two, so they
721 * reach the register file limit -- and start spilling to scratch -- well
722 * before a wider block can buy anything. Sweeping them by default was a
723 * regression. Set NEKO_EB_TUNE=1 to opt back into the sweep
724 */
725 if (getenv("NEKO_AUTOTUNE") ||
726 getenv("NEKO_EB") || !neko_eb_sweep()) {
727 const int c = neko_eb_env();
728 CASE_VECTOR_PADDED_SEL(LX, c);
729 sprintf(neko_log_buf, "Vector Ax : %d elem/block (untuned)",
730 NEKO_EB_SEL(LX, c));
731 log_message(neko_log_buf);
732 log_end_section();
733 return c;
734 }
735
736 HIP_CHECK(hipEventCreate(&start));
737 HIP_CHECK(hipEventCreate(&stop));
738
739 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
740 CASE_VECTOR_KSTEP_PADDED(LX, 0);
741 CASE_VECTOR_KSTEP_PADDED(LX, 1);
742 CASE_VECTOR_KSTEP_PADDED(LX, 2);
743 }
744
745 for (int r = 0; r < rounds; r++) {
746 NEKO_TUNE_TIME(time, CASE_VECTOR_KSTEP_PADDED, LX, 0, iters);
747 NEKO_TUNE_TIME(time, CASE_VECTOR_KSTEP_PADDED, LX, 1, iters);
748 NEKO_TUNE_TIME(time, CASE_VECTOR_KSTEP_PADDED, LX, 2, iters);
749 }
750
751 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
752 sprintf(neko_log_buf, "eb=%-4d : %9.2f us/call",
753 NEKO_EB_SEL(LX, c), time[c] * 10.0);
754 log_message(neko_log_buf);
755 }
756
757 NEKO_TUNE_BEST(time, best, NEKO_EB_CANDIDATES);
758
759 HIP_CHECK(hipEventDestroy(start));
760 HIP_CHECK(hipEventDestroy(stop));
761
762 /* Leave the output of the chosen kernel in au/av/aw */
763 CASE_VECTOR_PADDED_SEL(LX, best);
764
765 sprintf(neko_log_buf, "Chose : %d elem/block", NEKO_EB_SEL(LX, best));
766 log_message(neko_log_buf);
767 log_end_section();
768 return best;
769}
CASE_VECTOR
#define CASE_VECTOR(LX)
CASE_KSTEP_SEL
#define CASE_KSTEP_SEL(LX, SEL)
CASE
#define CASE(LX)
hip_ax_helm_vector_part2
void hip_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.hip:340
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.hip:358
CASE_VECTOR_KSTEP
#define CASE_VECTOR_KSTEP(LX, C)
CASE_VECTOR_PADDED_SEL
#define CASE_VECTOR_PADDED_SEL(LX, SEL)
CASE_PADDED
#define CASE_PADDED(LX)
CASE_KSTEP
#define CASE_KSTEP(LX, C)
CASE_KSTEP_PADDED_SEL
#define CASE_KSTEP_PADDED_SEL(LX, SEL)
CASE_LARGE
#define CASE_LARGE(LX)
tune_vector
int tune_vector(void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.hip:621
hip_ax_helm_vector
void hip_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.hip:236
tune_vector_padded
int tune_vector_padded(void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.hip:696
hip_ax_helm
void hip_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.hip:79
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.hip:487
CASE_VECTOR_KSTEP_PADDED
#define CASE_VECTOR_KSTEP_PADDED(LX, C)
CASE_1D
#define CASE_1D(LX, C)
CASE_VECTOR_PADDED
#define CASE_VECTOR_PADDED(LX)
CASE_VECTOR_SEL
#define CASE_VECTOR_SEL(LX, SEL)
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
HIP_CHECK
#define HIP_CHECK(err)
Definition check.h:8
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)