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