Neko 1.99.9
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 int *nwf_sel);
54
55template < const int>
56int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
57 void *dxt, void *dyt, void *dzt, void *h1,
58 void *g11, void *g22, void *g33, void *g12,
59 void *g13, void *g23, int *nelv, int *lx, int *eb_sel,
60 int *ch_sel, int *nwf_sel);
61
62template < const int>
63int tune_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
64 void *dx, void *dy, void *dz, void *h1,
65 void *g11, void *g22, void *g33, void *g12,
66 void *g13, void *g23, int *nelv, int *lx, int *eb_sel,
67 int *nwf_sel);
68
69template < const int>
70int tune_vector_padded(void *au, void *av, void *aw,
71 void *u, void *v, void *w,
72 void *dx, void *dy, void *dz, void *h1,
73 void *g11, void *g22, void *g33, void *g12,
74 void *g13, void *g23, int *nelv, int *lx, int *eb_sel,
75 int *nwf_sel);
76
77extern "C" {
78
82 void hip_ax_helm(void *w, void *u, void *dx, void *dy, void *dz,
83 void *dxt, void *dyt, void *dzt, void *h1,
84 void *g11, void *g22, void *g33, void *g12,
85 void *g13, void *g23, int *nelv, int *lx) {
86
87 static int autotune[13] = { 0 };
88 /* Elements per block candidate chosen by the autotuner, see
89 elem_block<> in elem_block.h */
90 static int autotune_eb[13] = { 0 };
91 /* chunk candidate chosen for the 1d variant */
92 static int autotune_ch[13] = { 0 };
93 /* wavefronts per block candidate chosen for the mfma variant */
94 static int autotune_nwf[13] = { 0 };
95
96 const dim3 nthrds_1d(1024, 1, 1);
97 const dim3 nblcks_1d((*nelv), 1, 1);
98
99#define CASE_1D(LX, C) \
100 hipLaunchKernelGGL(HIP_KERNEL_NAME( \
101 ax_helm_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> ), \
102 nblcks_1d, NEKO_CHUNKS_NTHRDS(LX, C), 0, \
103 (hipStream_t) glb_cmd_queue, \
104 (real *) w, (real *) u, \
105 (real *) dx, (real *) dy, (real *) dz, \
106 (real *) dxt, (real *) dyt, (real *) dzt, (real *) h1, \
107 (real *) g11, (real *) g22, (real *) g33, \
108 (real *) g12, (real *) g13, (real *) g23); \
109 HIP_CHECK(hipGetLastError());
110
111/* Runtime dispatch onto the tuned chunk candidate */
112#define CASE_1D_SEL(LX, SEL) \
113 switch (SEL) { \
114 case 0: CASE_1D(LX, 0); break; \
115 case 1: CASE_1D(LX, 1); break; \
116 case 2: CASE_1D(LX, 2); break; \
117 default: CASE_1D(LX, 3); break; \
118 }
119
120#define CASE_KSTEP(LX, C) \
121 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
122 ax_helm_kernel_kstep<real, LX, NEKO_EB(LX, C)> ), \
123 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
124 (hipStream_t) glb_cmd_queue, \
125 (real *) w, (real *) u, \
126 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
127 (real *) g11, (real *) g22, (real *) g33, \
128 (real *) g12, (real *) g13, (real *) g23, *nelv); \
129 HIP_CHECK(hipGetLastError());
130
131#define CASE_KSTEP_PADDED(LX, C) \
132 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
133 ax_helm_kernel_kstep_padded<real, LX, NEKO_EB(LX, C)> ),\
134 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
135 (hipStream_t) glb_cmd_queue, \
136 (real *) w, (real *) u, \
137 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
138 (real *) g11, (real *) g22, (real *) g33, \
139 (real *) g12, (real *) g13, (real *) g23, *nelv); \
140 HIP_CHECK(hipGetLastError());
141
142#define CASE_MFMA(LX, C) \
143 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
144 ax_helm_kernel_mfma<real, LX, NEKO_MFMA_NWF(C), \
145 NEKO_MFMA_TILE(C)> ), \
146 NEKO_MFMA_NBLCKS(*nelv, LX, C), NEKO_MFMA_NTHRDS(C), 0, \
147 (hipStream_t) glb_cmd_queue, \
148 (real *) w, (real *) u, \
149 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
150 (real *) g11, (real *) g22, (real *) g33, \
151 (real *) g12, (real *) g13, (real *) g23, *nelv); \
152 HIP_CHECK(hipGetLastError());
153
154/* Runtime dispatch onto the tuned candidate, which names both the wavefronts
155 per block and the matrix core tile, see NEKO_MFMA_TILE in mfma_kernel.h */
156#define CASE_MFMA_SEL(LX, SEL) \
157 switch (SEL) { \
158 case 0: CASE_MFMA(LX, 0); break; \
159 case 1: CASE_MFMA(LX, 1); break; \
160 case 2: CASE_MFMA(LX, 2); break; \
161 case 3: CASE_MFMA(LX, 3); break; \
162 case 4: CASE_MFMA(LX, 4); break; \
163 case 5: CASE_MFMA(LX, 5); break; \
164 case 6: CASE_MFMA(LX, 6); break; \
165 case 7: CASE_MFMA(LX, 7); break; \
166 case 8: CASE_MFMA(LX, 8); break; \
167 default: CASE_MFMA(LX, 9); break; \
168 }
169
170/* The switch above has to name every candidate. A candidate that fell
171 through to the default arm would launch its own grid and block shape --
172 NEKO_MFMA_NBLCKS and NEKO_MFMA_NTHRDS are evaluated at the call site --
173 at a kernel compiled for a different wavefront count, which is a launch
174 failure where the block is bigger and silent corruption where it is not.
175 Raising the candidate count without extending the switch stops here. */
176static_assert(NEKO_MFMA_CANDIDATES == 10,
177 "CASE_MFMA_SEL does not cover the candidate space");
178
179/* Runtime dispatch onto the tuned elements per block candidate */
180#define CASE_KSTEP_SEL(LX, SEL) \
181 switch (SEL) { \
182 case 0: CASE_KSTEP(LX, 0); break; \
183 case 1: CASE_KSTEP(LX, 1); break; \
184 default: CASE_KSTEP(LX, 2); break; \
185 }
186
187#define CASE_KSTEP_PADDED_SEL(LX, SEL) \
188 switch (SEL) { \
189 case 0: CASE_KSTEP_PADDED(LX, 0); break; \
190 case 1: CASE_KSTEP_PADDED(LX, 1); break; \
191 default: CASE_KSTEP_PADDED(LX, 2); break; \
192 }
193
194#define CASE(LX) \
195 case LX: \
196 if(autotune[LX] == 0 ) { \
197 autotune[LX]=tune<LX>( w, u, \
198 dx, dy, dz, \
199 dxt, dyt, dzt,h1, \
200 g11, g22, g33, \
201 g12, g13, g23, nelv, lx, \
202 &autotune_eb[LX], &autotune_ch[LX], \
203 &autotune_nwf[LX]); \
204 } else if (autotune[LX] == 1 ) { \
205 CASE_1D_SEL(LX, autotune_ch[LX]); \
206 } else if (autotune[LX] == 2 ) { \
207 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
208 } else if (autotune[LX] == 3 ) { \
209 CASE_MFMA_SEL(LX, autotune_nwf[LX]); \
210 } \
211 break
212
213#define CASE_PADDED(LX) \
214 case LX: \
215 if(autotune[LX] == 0 ) { \
216 autotune[LX]=tune_padded<LX>(w, u, \
217 dx, dy, dz, \
218 dxt, dyt, dzt,h1, \
219 g11, g22, g33, \
220 g12, g13, g23,nelv,lx, \
221 &autotune_eb[LX], &autotune_ch[LX], \
222 &autotune_nwf[LX]); \
223 } else if (autotune[LX] == 1 ) { \
224 CASE_1D_SEL(LX, autotune_ch[LX]); \
225 } else if (autotune[LX] == 2 ) { \
226 CASE_KSTEP_PADDED_SEL(LX, autotune_eb[LX]); \
227 } else if (autotune[LX] == 3 ) { \
228 CASE_MFMA_SEL(LX, autotune_nwf[LX]); \
229 } \
230 break
231
232/*
233 * High order cases have no 1d variant to compare against (its LDS footprint
234 * grows as LX^3), so they are not tuned and keep one element per block,
235 * i.e. candidate 0
236 */
237#define CASE_LARGE(LX) \
238 case LX: \
239 CASE_KSTEP(LX, 0); \
240 break
241
242#define CASE_LARGE_PADDED(LX) \
243 case LX: \
244 CASE_KSTEP_PADDED(LX, 0); \
245 break
246
247 if ((*lx) < 13) {
248 switch(*lx) {
249 CASE(2);
250 CASE(3);
251 CASE(4);
252 CASE(5);
253 CASE(6);
254 CASE(7);
255 CASE_PADDED(8);
256 CASE(9);
257 CASE(10);
258 CASE(11);
259 CASE(12);
260 default:
261 {
262 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
263 exit(1);
264 }
265 }
266 }
267 else {
268 switch(*lx) {
269 CASE_LARGE(13);
270 CASE_LARGE(14);
271 CASE_LARGE(15);
272 CASE_LARGE_PADDED(16);
273 default:
274 {
275 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
276 exit(1);
277 }
278 }
279 }
280 }
284 void hip_ax_helm_vector(void *au, void *av, void *aw,
285 void *u, void *v, void *w,
286 void *dx, void *dy, void *dz,
287 void *dxt, void *dyt, void *dzt,
288 void *h1, void *g11, void *g22,
289 void *g33, void *g12, void *g13,
290 void *g23, int *nelv, int *lx) {
291
292 /* Strategy chosen by the autotuner: 0 untuned, 2 kstep, 3 mfma. There
293 is no 1d variant of the vector operator, so the kstep family is the
294 baseline and the matrix core variant the only alternative */
295 static int autotune_v[17] = { 0 };
296 /* elements per block candidate chosen for the kstep variant */
297 static int autotune_v_eb[17] = { 0 };
298 /* wavefronts per block candidate chosen for the mfma variant */
299 static int autotune_v_nwf[17] = { 0 };
300
301#define CASE_VECTOR_KSTEP(LX, C) \
302 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
303 ax_helm_kernel_vector_kstep<real, LX, \
304 NEKO_EB(LX, C)> ), \
305 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
306 (hipStream_t) glb_cmd_queue, \
307 (real *) au, (real *) av, (real *) aw, \
308 (real *) u, (real *) v, (real *) w, \
309 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
310 (real *) g11, (real *) g22, (real *) g33, \
311 (real *) g12, (real *) g13, (real *) g23, *nelv); \
312 HIP_CHECK(hipGetLastError());
313
314#define CASE_VECTOR_KSTEP_PADDED(LX, C) \
315 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
316 ax_helm_kernel_vector_kstep_padded<real, LX, \
317 NEKO_EB(LX, C)> ), \
318 NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, \
319 (hipStream_t) glb_cmd_queue, \
320 (real *) au, (real *) av, (real *) aw, \
321 (real *) u, (real *) v, (real *) w, \
322 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
323 (real *) g11, (real *) g22, (real *) g33, \
324 (real *) g12, (real *) g13, (real *) g23, *nelv); \
325 HIP_CHECK(hipGetLastError());
326
327/* The matrix core variant: same grid and wavefronts per block candidates as
328 the scalar CASE_MFMA, see mfma_kernel.h */
329#define CASE_VECTOR_MFMA(LX, C) \
330 hipLaunchKernelGGL( HIP_KERNEL_NAME( \
331 ax_helm_kernel_mfma_vector<real, LX, \
332 NEKO_MFMA_NWF(C), \
333 NEKO_MFMA_TILE(C)> ), \
334 NEKO_MFMA_NBLCKS(*nelv, LX, C), NEKO_MFMA_NTHRDS(C), 0, \
335 (hipStream_t) glb_cmd_queue, \
336 (real *) au, (real *) av, (real *) aw, \
337 (real *) u, (real *) v, (real *) w, \
338 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
339 (real *) g11, (real *) g22, (real *) g33, \
340 (real *) g12, (real *) g13, (real *) g23, *nelv); \
341 HIP_CHECK(hipGetLastError());
342
343#define CASE_VECTOR_KSTEP_SEL(LX, SEL) \
344 switch (SEL) { \
345 case 0: CASE_VECTOR_KSTEP(LX, 0); break; \
346 case 1: CASE_VECTOR_KSTEP(LX, 1); break; \
347 default: CASE_VECTOR_KSTEP(LX, 2); break; \
348 }
349
350#define CASE_VECTOR_KSTEP_PADDED_SEL(LX, SEL) \
351 switch (SEL) { \
352 case 0: CASE_VECTOR_KSTEP_PADDED(LX, 0); break; \
353 case 1: CASE_VECTOR_KSTEP_PADDED(LX, 1); break; \
354 default: CASE_VECTOR_KSTEP_PADDED(LX, 2); break; \
355 }
356
357/* Runtime dispatch onto the tuned candidate, see CASE_MFMA_SEL */
358#define CASE_VECTOR_MFMA_SEL(LX, SEL) \
359 switch (SEL) { \
360 case 0: CASE_VECTOR_MFMA(LX, 0); break; \
361 case 1: CASE_VECTOR_MFMA(LX, 1); break; \
362 case 2: CASE_VECTOR_MFMA(LX, 2); break; \
363 case 3: CASE_VECTOR_MFMA(LX, 3); break; \
364 case 4: CASE_VECTOR_MFMA(LX, 4); break; \
365 case 5: CASE_VECTOR_MFMA(LX, 5); break; \
366 case 6: CASE_VECTOR_MFMA(LX, 6); break; \
367 case 7: CASE_VECTOR_MFMA(LX, 7); break; \
368 case 8: CASE_VECTOR_MFMA(LX, 8); break; \
369 default: CASE_VECTOR_MFMA(LX, 9); break; \
370 }
371
372/* The switch above has to name every candidate. A candidate that fell
373 through to the default arm would launch its own grid and block shape --
374 NEKO_MFMA_NBLCKS and NEKO_MFMA_NTHRDS are evaluated at the call site --
375 at a kernel compiled for a different wavefront count, which is a launch
376 failure where the block is bigger and silent corruption where it is not.
377 Raising the candidate count without extending the switch stops here. */
378static_assert(NEKO_MFMA_CANDIDATES == 10,
379 "CASE_VECTOR_MFMA_SEL does not cover the candidate space");
380
381#define CASE_VECTOR(LX) \
382 case LX: \
383 if (autotune_v[LX] == 0) { \
384 autotune_v[LX] = tune_vector<LX>(au, av, aw, u, v, w, \
385 dx, dy, dz, h1, \
386 g11, g22, g33, \
387 g12, g13, g23, nelv, lx, \
388 &autotune_v_eb[LX], \
389 &autotune_v_nwf[LX]); \
390 } else if (autotune_v[LX] == 2) { \
391 CASE_VECTOR_KSTEP_SEL(LX, autotune_v_eb[LX]); \
392 } else if (autotune_v[LX] == 3) { \
393 CASE_VECTOR_MFMA_SEL(LX, autotune_v_nwf[LX]); \
394 } \
395 break
396
397#define CASE_VECTOR_PADDED(LX) \
398 case LX: \
399 if (autotune_v[LX] == 0) { \
400 autotune_v[LX] = tune_vector_padded<LX>(au, av, aw, u, v, w, \
401 dx, dy, dz, h1, \
402 g11, g22, g33, \
403 g12, g13, g23, nelv, lx, \
404 &autotune_v_eb[LX], \
405 &autotune_v_nwf[LX]); \
406 } else if (autotune_v[LX] == 2) { \
407 CASE_VECTOR_KSTEP_PADDED_SEL(LX, autotune_v_eb[LX]); \
408 } else if (autotune_v[LX] == 3) { \
409 CASE_VECTOR_MFMA_SEL(LX, autotune_v_nwf[LX]); \
410 } \
411 break
412
413 switch(*lx) {
414 CASE_VECTOR(2);
415 CASE_VECTOR(3);
416 CASE_VECTOR_PADDED(4);
417 CASE_VECTOR(5);
418 CASE_VECTOR(6);
419 CASE_VECTOR(7);
420 CASE_VECTOR_PADDED(8);
421 CASE_VECTOR(9);
422 CASE_VECTOR(10);
423 CASE_VECTOR(11);
424 CASE_VECTOR(12);
425 CASE_VECTOR(13);
426 CASE_VECTOR(14);
427 CASE_VECTOR(15);
428 CASE_VECTOR_PADDED(16);
429 default:
430 {
431 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
432 exit(1);
433 }
434 }
435 }
436
440 void hip_ax_helm_vector_part2(void *au, void *av, void *aw,
441 void *u, void *v, void *w,
442 void *h2, void *B, int *n) {
443
444 const dim3 nthrds(1024, 1, 1);
445 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
446 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
447
448 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_vector_part2<real> ),
449 nblcks, nthrds, 0, stream,
450 (real *) au, (real *) av, (real *) aw,
451 (real *) u, (real *) v, (real *) w,
452 (real *) h2, (real *) B, *n);
453 }
454
455}
456
457template < const int LX >
458int tune(void *w, void *u, void *dx, void *dy, void *dz,
459 void *dxt, void *dyt, void *dzt, void *h1,
460 void *g11, void *g22, void *g33, void *g12,
461 void *g13, void *g23, int *nelv, int *lx, int *eb_sel, int *ch_sel,
462 int *nwf_sel) {
463 hipEvent_t start,stop;
464 float time1[NEKO_CHUNKS_CANDIDATES];
465 int best1 = 0;
466 float time2[NEKO_EB_CANDIDATES];
467 float time3[NEKO_MFMA_CANDIDATES];
468 int best3 = 0;
469 const int rounds = neko_tune_rounds();
470 const int iters = neko_tune_iters();
471 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
472 elements per block sweep is off */
473 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
474 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
475 const int mfma_cand = neko_mfma_candidates();
476 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
477 const int ch_pin = neko_chunks_pin();
478 const int eb_pin = neko_eb_pin();
479 const int nwf_pin = neko_mfma_pin();
480 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier tune() returns */
481 int strat = 0;
482 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
483 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
484 int best = 0;
485 int retval;
486
487 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
488 time2[c] = NEKO_TUNE_INIT;
489 }
490 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
491 time1[c] = NEKO_TUNE_INIT;
492 }
493 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
494 time3[c] = NEKO_TUNE_INIT;
495 }
496
497 const dim3 nthrds_1d(1024, 1, 1);
498 const dim3 nblcks_1d((*nelv), 1, 1);
499 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
500
501 char *env_value = NULL;
502 char neko_log_buf[80];
503
504 env_value=getenv("NEKO_AUTOTUNE");
505
506 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
507 log_section(neko_log_buf);
508
509 *eb_sel = 0;
510 *ch_sel = 0;
511 *nwf_sel = 0;
512
513 /*
514 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
515 * below is narrowed to that one kernel family, but its geometry -- the
516 * chunk size, the elements per block, the wavefronts per block -- is still
517 * measured candidate against candidate. A formulation this build or this
518 * device does not have is reported and ignored, leaving the full sweep.
519 */
520 if(env_value) {
521 if( !strcmp(env_value,"1D") ) {
522 strat = 1;
523 } else if( !strcmp(env_value,"KSTEP") ) {
524 strat = 2;
525 } else if( !strcmp(env_value,"MFMA") ) {
526 if (mfma) {
527 strat = 3;
528 } else {
529 sprintf(neko_log_buf, "MFMA strategy not available for this config");
530 log_error(neko_log_buf);
531 }
532 } else {
533 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
534 log_error(neko_log_buf);
535 }
536 }
537
538 /* Geometry of the pinned formulation, if its own variable fixes that too.
539 Both pinned leaves nothing to measure, so the kernel is launched once and
540 reported, which is what pinning has always done */
541 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
542 (strat == 3) ? nwf_pin : -1;
543
544 if (pin >= 0) {
545 switch (strat) {
546 case 1:
547 *ch_sel = pin;
548 CASE_1D_SEL(LX, pin);
549 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
550 NEKO_CHUNKS_SEL(LX, pin));
551 break;
552 case 2:
553 *eb_sel = pin;
554 CASE_KSTEP_SEL(LX, pin);
555 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
556 NEKO_EB_SEL(LX, pin));
557 break;
558 default:
559 *nwf_sel = pin;
560 CASE_MFMA_SEL(LX, pin);
561 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
562 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
563 break;
564 }
565 log_message(neko_log_buf);
566 log_end_section();
567 return strat;
568 }
569
570 if (strat) {
571 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
572 log_message(neko_log_buf);
573 }
574
575 /* Formulations the sweep considers, see NEKO_TUNE_FOR(). NEKO_MFMA_TUNE
576 keeps the matrix core variant out of an unpinned sweep, but naming it
577 explicitly still measures it */
578 const bool try_1d = (strat == 0 || strat == 1);
579 const bool try_kstep = (strat == 0 || strat == 2);
580 const bool try_mfma = mfma && ((strat == 3) ||
581 (strat == 0 && neko_mfma_sweep()));
582
583 HIP_CHECK(hipEventCreate(&start));
584 HIP_CHECK(hipEventCreate(&stop));
585
586 /* Warm every variant before timing anything: each specialisation has to
587 be resident and the clocks at steady state, or whichever variant is
588 timed first is measured on a colder part */
589 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
590 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
591 CASE_1D_SEL(LX, c);
592 }
593 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
594 CASE_KSTEP_SEL(LX, c);
595 }
596 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
597 CASE_MFMA_SEL(LX, c);
598 }
599 }
600
601 /* Interleaved rounds, best time per variant: timing them one after another
602 in a fixed order lets clock drift bias the comparison by position */
603 for (int r = 0; r < rounds; r++) {
604 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
605 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
606 }
607 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
608 NEKO_TUNE_TIME(time2, CASE_KSTEP_SEL, LX, c, iters);
609 }
610 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
611 NEKO_TUNE_TIME(time3, CASE_MFMA_SEL, LX, c, iters);
612 }
613 }
614
615 NEKO_TUNE_LOG(LX, time1, time2);
616 NEKO_TUNE_LOG_MFMA(LX, time3);
617
618 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
619 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
620 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
621 *eb_sel = best;
622 *ch_sel = best1;
623 *nwf_sel = best3;
624
625 if (time1[best1] < time2[best]) {
626 retval = 1;
627 } else {
628 retval = 2;
629 }
630
631 /* The mfma variant joins the comparison only where it exists, its
632 candidates are left at NEKO_TUNE_INIT otherwise */
633 if (time3[best3] < ((retval == 1) ? time1[best1] : time2[best])) {
634 retval = 3;
635 }
636
637 HIP_CHECK(hipEventDestroy(start));
638 HIP_CHECK(hipEventDestroy(stop));
639
640 /* The tuner stands in for a real Ax evaluation, and the variants do not
641 sum in the same order, so leave the output of the chosen kernel in w */
642 if (retval == 1) {
643 CASE_1D_SEL(LX, best1);
644 } else if (retval == 2) {
645 CASE_KSTEP_SEL(LX, best);
646 } else {
647 CASE_MFMA_SEL(LX, best3);
648 }
649
650 if (retval == 1) {
651 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
652 NEKO_CHUNKS_SEL(LX, best1));
653 } else if (retval == 2) {
654 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
655 NEKO_EB_SEL(LX, best));
656 } else {
657 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
658 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
659 NEKO_MFMA_EB(LX, best3));
660 }
661 log_message(neko_log_buf);
662 log_end_section();
663 return retval;
664}
665
666template < const int LX >
667int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
668 void *dxt, void *dyt, void *dzt, void *h1,
669 void *g11, void *g22, void *g33, void *g12,
670 void *g13, void *g23, int *nelv, int *lx,
671 int *eb_sel, int *ch_sel, int *nwf_sel) {
672 hipEvent_t start, stop;
673 float time1[NEKO_CHUNKS_CANDIDATES];
674 int best1 = 0;
675 float time2[NEKO_EB_CANDIDATES];
676 float time3[NEKO_MFMA_CANDIDATES];
677 int best3 = 0;
678 const int rounds = neko_tune_rounds();
679 const int iters = neko_tune_iters();
680 /* Candidates of the kstep sweep, see tune() */
681 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
682 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
683 const int mfma_cand = neko_mfma_candidates();
684 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
685 const int ch_pin = neko_chunks_pin();
686 const int eb_pin = neko_eb_pin();
687 const int nwf_pin = neko_mfma_pin();
688 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
689 int strat = 0;
690 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
691 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
692 int best = 0;
693 int retval;
694
695 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
696 time2[c] = NEKO_TUNE_INIT;
697 }
698 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
699 time1[c] = NEKO_TUNE_INIT;
700 }
701 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
702 time3[c] = NEKO_TUNE_INIT;
703 }
704
705 const dim3 nthrds_1d(1024, 1, 1);
706 const dim3 nblcks_1d((*nelv), 1, 1);
707 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
708
709 char *env_value = NULL;
710 char neko_log_buf[80];
711
712 env_value=getenv("NEKO_AUTOTUNE");
713
714 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
715 log_section(neko_log_buf);
716
717 *eb_sel = 0;
718 *ch_sel = 0;
719 *nwf_sel = 0;
720
721 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune() */
722 if(env_value) {
723 if( !strcmp(env_value,"1D") ) {
724 strat = 1;
725 } else if( !strcmp(env_value,"KSTEP") ) {
726 strat = 2;
727 } else if( !strcmp(env_value,"MFMA") ) {
728 if (mfma) {
729 strat = 3;
730 } else {
731 sprintf(neko_log_buf, "MFMA strategy not available for this config");
732 log_error(neko_log_buf);
733 }
734 } else {
735 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
736 log_error(neko_log_buf);
737 }
738 }
739
740 /* Formulation and geometry both pinned leaves nothing to measure */
741 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
742 (strat == 3) ? nwf_pin : -1;
743
744 if (pin >= 0) {
745 switch (strat) {
746 case 1:
747 *ch_sel = pin;
748 CASE_1D_SEL(LX, pin);
749 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
750 NEKO_CHUNKS_SEL(LX, pin));
751 break;
752 case 2:
753 *eb_sel = pin;
754 CASE_KSTEP_PADDED_SEL(LX, pin);
755 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
756 NEKO_EB_SEL(LX, pin));
757 break;
758 default:
759 *nwf_sel = pin;
760 CASE_MFMA_SEL(LX, pin);
761 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
762 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
763 break;
764 }
765 log_message(neko_log_buf);
766 log_end_section();
767 return strat;
768 }
769
770 if (strat) {
771 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
772 log_message(neko_log_buf);
773 }
774
775 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
776 const bool try_1d = (strat == 0 || strat == 1);
777 const bool try_kstep = (strat == 0 || strat == 2);
778 const bool try_mfma = mfma && ((strat == 3) ||
779 (strat == 0 && neko_mfma_sweep()));
780
781 HIP_CHECK(hipEventCreate(&start));
782 HIP_CHECK(hipEventCreate(&stop));
783
784 /* Warm every variant before timing anything: each specialisation has to
785 be resident and the clocks at steady state, or whichever variant is
786 timed first is measured on a colder part */
787 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
788 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
789 CASE_1D_SEL(LX, c);
790 }
791 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
792 CASE_KSTEP_PADDED_SEL(LX, c);
793 }
794 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
795 CASE_MFMA_SEL(LX, c);
796 }
797 }
798
799 /* Interleaved rounds, best time per variant: timing them one after another
800 in a fixed order lets clock drift bias the comparison by position */
801 for (int r = 0; r < rounds; r++) {
802 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
803 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
804 }
805 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
806 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED_SEL, LX, c, iters);
807 }
808 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
809 NEKO_TUNE_TIME(time3, CASE_MFMA_SEL, LX, c, iters);
810 }
811 }
812
813 NEKO_TUNE_LOG(LX, time1, time2);
814 NEKO_TUNE_LOG_MFMA(LX, time3);
815
816 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
817 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
818 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
819 *eb_sel = best;
820 *ch_sel = best1;
821 *nwf_sel = best3;
822
823 if (time1[best1] < time2[best]) {
824 retval = 1;
825 } else {
826 retval = 2;
827 }
828
829 /* The mfma variant joins the comparison only where it exists, its
830 candidates are left at NEKO_TUNE_INIT otherwise */
831 if (time3[best3] < ((retval == 1) ? time1[best1] : time2[best])) {
832 retval = 3;
833 }
834
835 HIP_CHECK(hipEventDestroy(start));
836 HIP_CHECK(hipEventDestroy(stop));
837
838 /* Leave the output of the chosen kernel in w, see tune() */
839 if (retval == 1) {
840 CASE_1D_SEL(LX, best1);
841 } else if (retval == 2) {
842 CASE_KSTEP_PADDED_SEL(LX, best);
843 } else {
844 CASE_MFMA_SEL(LX, best3);
845 }
846
847 if (retval == 1) {
848 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
849 NEKO_CHUNKS_SEL(LX, best1));
850 } else if (retval == 2) {
851 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
852 NEKO_EB_SEL(LX, best));
853 } else {
854 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
855 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
856 NEKO_MFMA_EB(LX, best3));
857 }
858 log_message(neko_log_buf);
859 log_end_section();
860 return retval;
861}
862
863/*
864 * The vector Ax has no 1d variant, so the tuner picks between the kstep
865 * family -- one element per block, plus the blocked candidates when the
866 * elements per block sweep is on -- and the matrix core variant, where the
867 * device and the polynomial order allow it. NEKO_AUTOTUNE=KSTEP (or 1D,
868 * there being nothing else to pin it to) narrows the sweep to the kstep
869 * variant, NEKO_AUTOTUNE=MFMA to the matrix core one, and in both cases the
870 * geometry candidates of the named variant are still measured -- it takes
871 * NEKO_EB or NEKO_MFMA_NWF to fix that too, as for the scalar operator.
872 *
873 * @note The section title has to fit the 30 character header field in
874 * log_section(), which truncates rather than wraps -- hence "Ax vector"
875 * rather than the "Ax helm vector" it reads as. It matches the CUDA copy.
876 */
877template < const int LX >
878int tune_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
879 void *dx, void *dy, void *dz, void *h1,
880 void *g11, void *g22, void *g33, void *g12,
881 void *g13, void *g23, int *nelv, int *lx, int *eb_sel,
882 int *nwf_sel) {
883 hipEvent_t start, stop;
884 float time2[NEKO_EB_CANDIDATES];
885 int best2 = 0;
886 float time3[NEKO_MFMA_CANDIDATES];
887 int best3 = 0;
888 const int rounds = neko_tune_rounds();
889 const int iters = neko_tune_iters();
890 const int sweep = neko_eb_sweep();
891 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
892 elements per block sweep is off */
893 const int eb_cand = sweep ? NEKO_EB_CANDIDATES : 1;
894 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
895 const int mfma_cand = neko_mfma_candidates();
896 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
897 const int eb_pin = neko_eb_pin();
898 const int nwf_pin = neko_mfma_pin();
899 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
900 int strat = 0;
901 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
902 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
903 int retval;
904
905 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
906 time2[c] = NEKO_TUNE_INIT;
907 }
908 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
909 time3[c] = NEKO_TUNE_INIT;
910 }
911
912 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
913
914 char *env_value = NULL;
915 char neko_log_buf[80];
916
917 env_value = getenv("NEKO_AUTOTUNE");
918
919 sprintf(neko_log_buf, "Autotune Ax vector (lx: %d)", *lx);
920 log_section(neko_log_buf);
921
922 *eb_sel = 0;
923 *nwf_sel = 0;
924
925 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune() */
926 if (env_value) {
927 if ( !strcmp(env_value, "KSTEP") || !strcmp(env_value, "1D") ) {
928 /* No 1d variant here, so either pin lands on kstep */
929 strat = 2;
930 } else if ( !strcmp(env_value, "MFMA") ) {
931 if (mfma) {
932 strat = 3;
933 } else {
934 sprintf(neko_log_buf, "MFMA strategy not available for this config");
935 log_error(neko_log_buf);
936 }
937 } else {
938 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
939 log_error(neko_log_buf);
940 }
941 }
942
943 /* Formulation and geometry both pinned leaves nothing to measure */
944 const int pin = (strat == 2) ? eb_pin : (strat == 3) ? nwf_pin : -1;
945
946 if (pin >= 0) {
947 if (strat == 2) {
948 *eb_sel = pin;
949 CASE_VECTOR_KSTEP_SEL(LX, pin);
950 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
951 NEKO_EB_SEL(LX, pin));
952 } else {
953 *nwf_sel = pin;
954 CASE_VECTOR_MFMA_SEL(LX, pin);
955 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
956 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
957 }
958 log_message(neko_log_buf);
959 log_end_section();
960 return strat;
961 }
962
963 if (strat) {
964 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
965 log_message(neko_log_buf);
966 }
967
968 /* Formulations the sweep considers, see NEKO_TUNE_FOR(). NEKO_MFMA_TUNE
969 keeps the matrix core variant out of an unpinned sweep, but naming it
970 explicitly still measures it */
971 const bool try_kstep = (strat == 0 || strat == 2);
972 const bool try_mfma = mfma && ((strat == 3) ||
973 (strat == 0 && neko_mfma_sweep()));
974
975 /* With neither the blocking sweep nor a matrix core variant there is one
976 candidate, which is what these kernels have always run */
977 if (!sweep && !try_mfma) {
978 CASE_VECTOR_KSTEP(LX, 0);
979 sprintf(neko_log_buf, "Vector Ax : %d elem/block (untuned)",
980 NEKO_EB_SEL(LX, 0));
981 log_message(neko_log_buf);
982 log_end_section();
983 return 2;
984 }
985
986 HIP_CHECK(hipEventCreate(&start));
987 HIP_CHECK(hipEventCreate(&stop));
988
989 /* Warm every variant before timing anything, see tune() */
990 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
991 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
992 CASE_VECTOR_KSTEP_SEL(LX, c);
993 }
994 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
995 CASE_VECTOR_MFMA_SEL(LX, c);
996 }
997 }
998
999 /* Interleaved rounds, best time per variant, see tune() */
1000 for (int r = 0; r < rounds; r++) {
1001 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1002 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_SEL, LX, c, iters);
1003 }
1004 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1005 NEKO_TUNE_TIME(time3, CASE_VECTOR_MFMA_SEL, LX, c, iters);
1006 }
1007 }
1008
1009 NEKO_TUNE_LOG_KSTEP(LX, time2);
1010 NEKO_TUNE_LOG_MFMA_VEC(LX, time3);
1011
1012 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1013 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
1014 *eb_sel = best2;
1015 *nwf_sel = best3;
1016
1017 /* The mfma variant joins the comparison only where it exists, its
1018 candidates are left at NEKO_TUNE_INIT otherwise */
1019 retval = (time3[best3] < time2[best2]) ? 3 : 2;
1020
1021 HIP_CHECK(hipEventDestroy(start));
1022 HIP_CHECK(hipEventDestroy(stop));
1023
1024 /* The tuner stands in for a real Ax evaluation, and the variants do not
1025 sum in the same order, so leave the output of the chosen kernel in
1026 au, av and aw */
1027 if (retval == 2) {
1028 CASE_VECTOR_KSTEP_SEL(LX, best2);
1029 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1030 NEKO_EB_SEL(LX, best2));
1031 } else {
1032 CASE_VECTOR_MFMA_SEL(LX, best3);
1033 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
1034 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
1035 NEKO_MFMA_EB(LX, best3));
1036 }
1037 log_message(neko_log_buf);
1038 log_end_section();
1039 return retval;
1040}
1041
1042/* Padded variant of tune_vector(), for the orders whose kstep kernel stages
1043 with a padded stride. The matrix core variant is the same either way */
1044template < const int LX >
1045int tune_vector_padded(void *au, void *av, void *aw,
1046 void *u, void *v, void *w,
1047 void *dx, void *dy, void *dz, void *h1,
1048 void *g11, void *g22, void *g33, void *g12,
1049 void *g13, void *g23, int *nelv, int *lx, int *eb_sel,
1050 int *nwf_sel) {
1051 hipEvent_t start, stop;
1052 float time2[NEKO_EB_CANDIDATES];
1053 int best2 = 0;
1054 float time3[NEKO_MFMA_CANDIDATES];
1055 int best3 = 0;
1056 const int rounds = neko_tune_rounds();
1057 const int iters = neko_tune_iters();
1058 const int sweep = neko_eb_sweep();
1059 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
1060 elements per block sweep is off */
1061 const int eb_cand = sweep ? NEKO_EB_CANDIDATES : 1;
1062 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
1063 const int mfma_cand = neko_mfma_candidates();
1064 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
1065 const int eb_pin = neko_eb_pin();
1066 const int nwf_pin = neko_mfma_pin();
1067 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
1068 int strat = 0;
1069 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
1070 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
1071 int retval;
1072
1073 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1074 time2[c] = NEKO_TUNE_INIT;
1075 }
1076 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
1077 time3[c] = NEKO_TUNE_INIT;
1078 }
1079
1080 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
1081
1082 char *env_value = NULL;
1083 char neko_log_buf[80];
1084
1085 env_value = getenv("NEKO_AUTOTUNE");
1086
1087 sprintf(neko_log_buf, "Autotune Ax vector (lx: %d)", *lx);
1088 log_section(neko_log_buf);
1089
1090 *eb_sel = 0;
1091 *nwf_sel = 0;
1092
1093 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune() */
1094 if (env_value) {
1095 if ( !strcmp(env_value, "KSTEP") || !strcmp(env_value, "1D") ) {
1096 /* No 1d variant here, so either pin lands on kstep */
1097 strat = 2;
1098 } else if ( !strcmp(env_value, "MFMA") ) {
1099 if (mfma) {
1100 strat = 3;
1101 } else {
1102 sprintf(neko_log_buf, "MFMA strategy not available for this config");
1103 log_error(neko_log_buf);
1104 }
1105 } else {
1106 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1107 log_error(neko_log_buf);
1108 }
1109 }
1110
1111 /* Formulation and geometry both pinned leaves nothing to measure */
1112 const int pin = (strat == 2) ? eb_pin : (strat == 3) ? nwf_pin : -1;
1113
1114 if (pin >= 0) {
1115 if (strat == 2) {
1116 *eb_sel = pin;
1117 CASE_VECTOR_KSTEP_PADDED_SEL(LX, pin);
1118 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
1119 NEKO_EB_SEL(LX, pin));
1120 } else {
1121 *nwf_sel = pin;
1122 CASE_VECTOR_MFMA_SEL(LX, pin);
1123 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
1124 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
1125 }
1126 log_message(neko_log_buf);
1127 log_end_section();
1128 return strat;
1129 }
1130
1131 if (strat) {
1132 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
1133 log_message(neko_log_buf);
1134 }
1135
1136 /* Formulations the sweep considers, see NEKO_TUNE_FOR(). NEKO_MFMA_TUNE
1137 keeps the matrix core variant out of an unpinned sweep, but naming it
1138 explicitly still measures it */
1139 const bool try_kstep = (strat == 0 || strat == 2);
1140 const bool try_mfma = mfma && ((strat == 3) ||
1141 (strat == 0 && neko_mfma_sweep()));
1142
1143 /* One candidate, see tune_vector() */
1144 if (!sweep && !try_mfma) {
1145 CASE_VECTOR_KSTEP_PADDED(LX, 0);
1146 sprintf(neko_log_buf, "Vector Ax : %d elem/block (untuned)",
1147 NEKO_EB_SEL(LX, 0));
1148 log_message(neko_log_buf);
1149 log_end_section();
1150 return 2;
1151 }
1152
1153 HIP_CHECK(hipEventCreate(&start));
1154 HIP_CHECK(hipEventCreate(&stop));
1155
1156 /* Warm every variant before timing anything, see tune() */
1157 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1158 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1159 CASE_VECTOR_KSTEP_PADDED_SEL(LX, c);
1160 }
1161 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1162 CASE_VECTOR_MFMA_SEL(LX, c);
1163 }
1164 }
1165
1166 /* Interleaved rounds, best time per variant, see tune() */
1167 for (int r = 0; r < rounds; r++) {
1168 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1169 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_PADDED_SEL, LX, c, iters);
1170 }
1171 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1172 NEKO_TUNE_TIME(time3, CASE_VECTOR_MFMA_SEL, LX, c, iters);
1173 }
1174 }
1175
1176 NEKO_TUNE_LOG_KSTEP(LX, time2);
1177 NEKO_TUNE_LOG_MFMA_VEC(LX, time3);
1178
1179 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1180 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
1181 *eb_sel = best2;
1182 *nwf_sel = best3;
1183
1184 /* The mfma variant joins the comparison only where it exists, its
1185 candidates are left at NEKO_TUNE_INIT otherwise */
1186 retval = (time3[best3] < time2[best2]) ? 3 : 2;
1187
1188 HIP_CHECK(hipEventDestroy(start));
1189 HIP_CHECK(hipEventDestroy(stop));
1190
1191 /* Leave the output of the chosen kernel in au, av and aw, see
1192 tune_vector() */
1193 if (retval == 2) {
1194 CASE_VECTOR_KSTEP_PADDED_SEL(LX, best2);
1195 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1196 NEKO_EB_SEL(LX, best2));
1197 } else {
1198 CASE_VECTOR_MFMA_SEL(LX, best3);
1199 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
1200 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
1201 NEKO_MFMA_EB(LX, best3));
1202 }
1203 log_message(neko_log_buf);
1204 log_end_section();
1205 return retval;
1206}
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, int *nwf_sel)
Definition ax_helm.hip:667
CASE_VECTOR
#define CASE_VECTOR(LX)
CASE_VECTOR_MFMA_SEL
#define CASE_VECTOR_MFMA_SEL(LX, SEL)
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:440
CASE_1D_SEL
#define CASE_1D_SEL(LX, SEL)
CASE_VECTOR_KSTEP
#define CASE_VECTOR_KSTEP(LX, C)
CASE_MFMA_SEL
#define CASE_MFMA_SEL(LX, SEL)
CASE_PADDED
#define CASE_PADDED(LX)
CASE_KSTEP_PADDED_SEL
#define CASE_KSTEP_PADDED_SEL(LX, SEL)
CASE_LARGE
#define CASE_LARGE(LX)
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:284
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:82
CASE_LARGE_PADDED
#define CASE_LARGE_PADDED(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, int *eb_sel, int *nwf_sel)
Definition ax_helm.hip:878
CASE_VECTOR_KSTEP_PADDED
#define CASE_VECTOR_KSTEP_PADDED(LX, C)
CASE_VECTOR_KSTEP_SEL
#define CASE_VECTOR_KSTEP_SEL(LX, SEL)
CASE_VECTOR_PADDED
#define CASE_VECTOR_PADDED(LX)
CASE_VECTOR_KSTEP_PADDED_SEL
#define CASE_VECTOR_KSTEP_PADDED_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, int *nwf_sel)
Definition ax_helm.hip:458
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, int *eb_sel, int *nwf_sel)
Definition ax_helm.hip:1045
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
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
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
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
av
__global__ void T *__restrict__ av
Definition ax_helm_full_kernel.h:44
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
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
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
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:184
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:180
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:183
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:185
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:182
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:181
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:179
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:115
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:116
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition cdtp_kernel.h:114
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_chunks_pin
static int neko_chunks_pin()
Definition elem_block_tune.h:112
NEKO_TUNE_TIME
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
Definition elem_block_tune.h:163
neko_tune_rounds
static int neko_tune_rounds()
Definition elem_block_tune.h:128
neko_eb_pin
static int neko_eb_pin()
Definition elem_block_tune.h:95
NEKO_TUNE_LOG
#define NEKO_TUNE_LOG(LX, T1, T2)
Definition elem_block_tune.h:193
NEKO_TUNE_INIT
#define NEKO_TUNE_INIT
Definition elem_block_tune.h:72
NEKO_TUNE_BEST
#define NEKO_TUNE_BEST(T, BEST, N)
Definition elem_block_tune.h:185
neko_tune_iters
static int neko_tune_iters()
Definition elem_block_tune.h:136
NEKO_TUNE_FOR
#define NEKO_TUNE_FOR(C, ON, PIN, N)
Definition elem_block_tune.h:157
neko_eb_sweep
static int neko_eb_sweep()
Definition elem_block_tune.h:75
NEKO_TUNE_WARMUP
#define NEKO_TUNE_WARMUP
Definition elem_block_tune.h:73
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_TUNE_LOG_KSTEP
#define NEKO_TUNE_LOG_KSTEP(LX, T2)
Definition elem_block_tune.h:210
NEKO_TUNE_LOG_MFMA_VEC
#define NEKO_TUNE_LOG_MFMA_VEC(LX, T3)
Definition mfma_kernel.h:623
NEKO_TUNE_LOG_MFMA
#define NEKO_TUNE_LOG_MFMA(LX, T3)
Definition mfma_kernel.h:602
mfma_tile_name
static const char * mfma_tile_name(const int c)
Definition mfma_kernel.h:374
NEKO_MFMA_CANDIDATES
#define NEKO_MFMA_CANDIDATES
Definition mfma_kernel.h:261
neko_mfma_pin
static int neko_mfma_pin()
Definition mfma_kernel.h:533
neko_mfma_candidates
static int neko_mfma_candidates()
Definition mfma_kernel.h:593
hip_have_mfma
static bool hip_have_mfma()
Definition mfma_kernel.h:136
NEKO_MFMA_NWF
#define NEKO_MFMA_NWF(C)
Definition mfma_kernel.h:263
NEKO_MFMA_EB
#define NEKO_MFMA_EB(LX, C)
Definition mfma_kernel.h:442
neko_mfma_sweep
static int neko_mfma_sweep()
Definition mfma_kernel.h:511
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)