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 /* Verify every candidate against the reference formulation on a synthetic
592 input before anything is timed, see neko_tune_verify_begin() in
593 elem_block_tune.h. A candidate that does not reproduce it is reported
594 and dropped from the comparison below. */
595 neko_tune_rej_t rej1 = {{0}}, rej2 = {{0}}, rej3 = {{0}};
596 {
597 neko_tune_verify_t vf;
598 real *vouts[] = { (real *) w };
599 void **vins[] = { &u };
600 if (neko_tune_verify_begin(&vf, vouts, 1, vins, 1,
601 (*nelv) * LX * LX * LX, stream)) {
602 sprintf(neko_log_buf, "Verify : rel tol %.1e, synthetic input",
603 vf.tol);
604 log_message(neko_log_buf);
605 CASE_1D_SEL(LX, 0);
606 neko_tune_verify_ref(&vf, stream);
607 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
608 NEKO_TUNE_VERIFY_CAND(vf, rej1, CASE_1D_SEL, LX, c, "1D");
609 }
610 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
611 NEKO_TUNE_VERIFY_CAND(vf, rej2, CASE_KSTEP_SEL, LX, c, "KSTEP");
612 }
613 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
614 NEKO_TUNE_VERIFY_CAND(vf, rej3, CASE_MFMA_SEL, LX, c, "MFMA");
615 }
616 neko_tune_verify_end(&vf, stream);
617 } else {
618 sprintf(neko_log_buf, "Verify : %s", neko_tune_verify_on() ?
619 "skipped, no memory for it" : "off (NEKO_TUNE_VERIFY=0)");
620 log_message(neko_log_buf);
621 }
622 }
623
624 /* Warm every variant before timing anything: each specialisation has to
625 be resident and the clocks at steady state, or whichever variant is
626 timed first is measured on a colder part */
627 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
628 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
629 CASE_1D_SEL(LX, c);
630 }
631 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
632 CASE_KSTEP_SEL(LX, c);
633 }
634 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
635 CASE_MFMA_SEL(LX, c);
636 }
637 }
638
639 /* Interleaved rounds, best time per variant: timing them one after another
640 in a fixed order lets clock drift bias the comparison by position */
641 for (int r = 0; r < rounds; r++) {
642 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
643 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
644 }
645 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
646 NEKO_TUNE_TIME(time2, CASE_KSTEP_SEL, LX, c, iters);
647 }
648 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
649 NEKO_TUNE_TIME(time3, CASE_MFMA_SEL, LX, c, iters);
650 }
651 }
652
653 /* Candidates the verification pass rejected drop out here, after the
654 timing, so the sweep's own bookkeeping is untouched */
655 NEKO_TUNE_APPLY_REJ(time1, rej1, NEKO_CHUNKS_CANDIDATES);
656 NEKO_TUNE_APPLY_REJ(time2, rej2, NEKO_EB_CANDIDATES);
657 NEKO_TUNE_APPLY_REJ(time3, rej3, NEKO_MFMA_CANDIDATES);
658
659 NEKO_TUNE_LOG(LX, time1, time2);
660 NEKO_TUNE_LOG_MFMA(LX, time3);
661
662 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
663 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
664 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
665 *eb_sel = best;
666 *ch_sel = best1;
667 *nwf_sel = best3;
668
669 if (time1[best1] < time2[best]) {
670 retval = 1;
671 } else {
672 retval = 2;
673 }
674
675 /* The mfma variant joins the comparison only where it exists, its
676 candidates are left at NEKO_TUNE_INIT otherwise */
677 if (time3[best3] < ((retval == 1) ? time1[best1] : time2[best])) {
678 retval = 3;
679 }
680
681 HIP_CHECK(hipEventDestroy(start));
682 HIP_CHECK(hipEventDestroy(stop));
683
684 /* The tuner stands in for a real Ax evaluation, and the variants do not
685 sum in the same order, so leave the output of the chosen kernel in w */
686 if (retval == 1) {
687 CASE_1D_SEL(LX, best1);
688 } else if (retval == 2) {
689 CASE_KSTEP_SEL(LX, best);
690 } else {
691 CASE_MFMA_SEL(LX, best3);
692 }
693
694 if (retval == 1) {
695 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
696 NEKO_CHUNKS_SEL(LX, best1));
697 } else if (retval == 2) {
698 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
699 NEKO_EB_SEL(LX, best));
700 } else {
701 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
702 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
703 NEKO_MFMA_EB(LX, best3));
704 }
705 log_message(neko_log_buf);
706 log_end_section();
707 return retval;
708}
709
710template < const int LX >
711int tune_compr_padded(void *w, void *u, void *dx, void *dy, void *dz,
712 void *dxt, void *dyt, void *dzt, void *h1,
713 void *g11, void *g22, void *g33, void *g12,
714 void *g13, void *g23, void *ind, int *nelv, int *lx,
715 int *eb_sel, int *ch_sel, int *nwf_sel) {
716 hipEvent_t start, stop;
717 float time1[NEKO_CHUNKS_CANDIDATES];
718 int best1 = 0;
719 float time2[NEKO_EB_CANDIDATES];
720 float time3[NEKO_MFMA_CANDIDATES];
721 int best3 = 0;
722 const int rounds = neko_tune_rounds();
723 const int iters = neko_tune_iters();
724 /* Candidates of the kstep sweep, see tune_compr() */
725 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
726 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
727 const int mfma_cand = neko_mfma_candidates();
728 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
729 const int ch_pin = neko_chunks_pin();
730 const int eb_pin = neko_eb_pin();
731 const int nwf_pin = neko_mfma_pin();
732 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
733 int strat = 0;
734 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
735 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
736 int best = 0;
737 int retval;
738
739 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
740 time2[c] = NEKO_TUNE_INIT;
741 }
742 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
743 time1[c] = NEKO_TUNE_INIT;
744 }
745 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
746 time3[c] = NEKO_TUNE_INIT;
747 }
748
749 const dim3 nthrds_1d(1024, 1, 1);
750 const dim3 nblcks_1d((*nelv), 1, 1);
751 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
752
753 char *env_value = NULL;
754 char neko_log_buf[80];
755
756 env_value=getenv("NEKO_AUTOTUNE");
757
758 sprintf(neko_log_buf, "Autotune Ax compr (lx: %d)", *lx);
759 log_section(neko_log_buf);
760
761 *eb_sel = 0;
762 *ch_sel = 0;
763 *nwf_sel = 0;
764
765 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune_compr() */
766 if(env_value) {
767 if( !strcmp(env_value,"1D") ) {
768 strat = 1;
769 } else if( !strcmp(env_value,"KSTEP") ) {
770 strat = 2;
771 } else if( !strcmp(env_value,"MFMA") ) {
772 if (mfma) {
773 strat = 3;
774 } else {
775 sprintf(neko_log_buf, "MFMA strategy not available for this config");
776 log_error(neko_log_buf);
777 }
778 } else {
779 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
780 log_error(neko_log_buf);
781 }
782 }
783
784 /* Formulation and geometry both pinned leaves nothing to measure */
785 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
786 (strat == 3) ? nwf_pin : -1;
787
788 if (pin >= 0) {
789 switch (strat) {
790 case 1:
791 *ch_sel = pin;
792 CASE_1D_SEL(LX, pin);
793 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
794 NEKO_CHUNKS_SEL(LX, pin));
795 break;
796 case 2:
797 *eb_sel = pin;
798 CASE_KSTEP_PADDED_SEL(LX, pin);
799 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
800 NEKO_EB_SEL(LX, pin));
801 break;
802 default:
803 *nwf_sel = pin;
804 CASE_MFMA_SEL(LX, pin);
805 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
806 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
807 break;
808 }
809 log_message(neko_log_buf);
810 log_end_section();
811 return strat;
812 }
813
814 if (strat) {
815 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
816 log_message(neko_log_buf);
817 }
818
819 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
820 const bool try_1d = (strat == 0 || strat == 1);
821 const bool try_kstep = (strat == 0 || strat == 2);
822 const bool try_mfma = mfma && ((strat == 3) ||
823 (strat == 0 && neko_mfma_sweep()));
824
825 HIP_CHECK(hipEventCreate(&start));
826 HIP_CHECK(hipEventCreate(&stop));
827
828 /* Verify every candidate against the reference formulation on a synthetic
829 input before anything is timed, see neko_tune_verify_begin() in
830 elem_block_tune.h. A candidate that does not reproduce it is reported
831 and dropped from the comparison below. */
832 neko_tune_rej_t rej1 = {{0}}, rej2 = {{0}}, rej3 = {{0}};
833 {
834 neko_tune_verify_t vf;
835 real *vouts[] = { (real *) w };
836 void **vins[] = { &u };
837 if (neko_tune_verify_begin(&vf, vouts, 1, vins, 1,
838 (*nelv) * LX * LX * LX, stream)) {
839 sprintf(neko_log_buf, "Verify : rel tol %.1e, synthetic input",
840 vf.tol);
841 log_message(neko_log_buf);
842 CASE_1D_SEL(LX, 0);
843 neko_tune_verify_ref(&vf, stream);
844 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
845 NEKO_TUNE_VERIFY_CAND(vf, rej1, CASE_1D_SEL, LX, c, "1D");
846 }
847 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
848 NEKO_TUNE_VERIFY_CAND(vf, rej2, CASE_KSTEP_PADDED_SEL, LX, c, "KSTEP");
849 }
850 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
851 NEKO_TUNE_VERIFY_CAND(vf, rej3, CASE_MFMA_SEL, LX, c, "MFMA");
852 }
853 neko_tune_verify_end(&vf, stream);
854 } else {
855 sprintf(neko_log_buf, "Verify : %s", neko_tune_verify_on() ?
856 "skipped, no memory for it" : "off (NEKO_TUNE_VERIFY=0)");
857 log_message(neko_log_buf);
858 }
859 }
860
861 /* Warm every variant before timing anything: each specialisation has to
862 be resident and the clocks at steady state, or whichever variant is
863 timed first is measured on a colder part */
864 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
865 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
866 CASE_1D_SEL(LX, c);
867 }
868 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
869 CASE_KSTEP_PADDED_SEL(LX, c);
870 }
871 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
872 CASE_MFMA_SEL(LX, c);
873 }
874 }
875
876 /* Interleaved rounds, best time per variant: timing them one after another
877 in a fixed order lets clock drift bias the comparison by position */
878 for (int r = 0; r < rounds; r++) {
879 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
880 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
881 }
882 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
883 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED_SEL, LX, c, iters);
884 }
885 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
886 NEKO_TUNE_TIME(time3, CASE_MFMA_SEL, LX, c, iters);
887 }
888 }
889
890 /* Candidates the verification pass rejected drop out here, after the
891 timing, so the sweep's own bookkeeping is untouched */
892 NEKO_TUNE_APPLY_REJ(time1, rej1, NEKO_CHUNKS_CANDIDATES);
893 NEKO_TUNE_APPLY_REJ(time2, rej2, NEKO_EB_CANDIDATES);
894 NEKO_TUNE_APPLY_REJ(time3, rej3, NEKO_MFMA_CANDIDATES);
895
896 NEKO_TUNE_LOG(LX, time1, time2);
897 NEKO_TUNE_LOG_MFMA(LX, time3);
898
899 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
900 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
901 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
902 *eb_sel = best;
903 *ch_sel = best1;
904 *nwf_sel = best3;
905
906 if (time1[best1] < time2[best]) {
907 retval = 1;
908 } else {
909 retval = 2;
910 }
911
912 /* The mfma variant joins the comparison only where it exists, its
913 candidates are left at NEKO_TUNE_INIT otherwise */
914 if (time3[best3] < ((retval == 1) ? time1[best1] : time2[best])) {
915 retval = 3;
916 }
917
918 HIP_CHECK(hipEventDestroy(start));
919 HIP_CHECK(hipEventDestroy(stop));
920
921 /* Leave the output of the chosen kernel in w, see tune_compr() */
922 if (retval == 1) {
923 CASE_1D_SEL(LX, best1);
924 } else if (retval == 2) {
925 CASE_KSTEP_PADDED_SEL(LX, best);
926 } else {
927 CASE_MFMA_SEL(LX, best3);
928 }
929
930 if (retval == 1) {
931 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
932 NEKO_CHUNKS_SEL(LX, best1));
933 } else if (retval == 2) {
934 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
935 NEKO_EB_SEL(LX, best));
936 } else {
937 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
938 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
939 NEKO_MFMA_EB(LX, best3));
940 }
941 log_message(neko_log_buf);
942 log_end_section();
943 return retval;
944}
945
946/*
947 * The vector Ax has no 1d variant, so the tuner picks between the kstep
948 * family -- one element per block, plus the blocked candidates when the
949 * elements per block sweep is on -- and the matrix core variant, where the
950 * device and the polynomial order allow it. NEKO_AUTOTUNE=KSTEP (or 1D,
951 * there being nothing else to pin it to) narrows the sweep to the kstep
952 * variant, NEKO_AUTOTUNE=MFMA to the matrix core one, and in both cases the
953 * geometry candidates of the named variant are still measured -- it takes
954 * NEKO_EB or NEKO_MFMA_NWF to fix that too, as for the scalar operator.
955 *
956 * @note The section title has to fit the 30 character header field in
957 * log_section(), which truncates rather than wraps -- hence "Ax compr vec"
958 * rather than the "Ax helm compr vector" it reads as, which fills the field
959 * exactly at a two digit lx.
960 */
961template < const int LX >
962int tune_compr_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
963 void *dx, void *dy, void *dz, void *h1,
964 void *g11, void *g22, void *g33, void *g12,
965 void *g13, void *g23, void *ind, int *nelv, int *lx,
966 int *eb_sel, int *nwf_sel) {
967 hipEvent_t start, stop;
968 float time2[NEKO_EB_CANDIDATES];
969 int best2 = 0;
970 float time3[NEKO_MFMA_CANDIDATES];
971 int best3 = 0;
972 const int rounds = neko_tune_rounds();
973 const int iters = neko_tune_iters();
974 const int sweep = neko_eb_sweep();
975 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
976 elements per block sweep is off */
977 const int eb_cand = sweep ? NEKO_EB_CANDIDATES : 1;
978 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
979 const int mfma_cand = neko_mfma_candidates();
980 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
981 const int eb_pin = neko_eb_pin();
982 const int nwf_pin = neko_mfma_pin();
983 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
984 int strat = 0;
985 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
986 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
987 int retval;
988
989 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
990 time2[c] = NEKO_TUNE_INIT;
991 }
992 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
993 time3[c] = NEKO_TUNE_INIT;
994 }
995
996 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
997
998 char *env_value = NULL;
999 char neko_log_buf[80];
1000
1001 env_value = getenv("NEKO_AUTOTUNE");
1002
1003 sprintf(neko_log_buf, "Autotune Ax compr vec (lx: %d)", *lx);
1004 log_section(neko_log_buf);
1005
1006 *eb_sel = 0;
1007 *nwf_sel = 0;
1008
1009 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune_compr() */
1010 if (env_value) {
1011 if ( !strcmp(env_value, "KSTEP") || !strcmp(env_value, "1D") ) {
1012 /* No 1d variant here, so either pin lands on kstep */
1013 strat = 2;
1014 } else if ( !strcmp(env_value, "MFMA") ) {
1015 if (mfma) {
1016 strat = 3;
1017 } else {
1018 sprintf(neko_log_buf, "MFMA strategy not available for this config");
1019 log_error(neko_log_buf);
1020 }
1021 } else {
1022 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1023 log_error(neko_log_buf);
1024 }
1025 }
1026
1027 /* Formulation and geometry both pinned leaves nothing to measure */
1028 const int pin = (strat == 2) ? eb_pin : (strat == 3) ? nwf_pin : -1;
1029
1030 if (pin >= 0) {
1031 if (strat == 2) {
1032 *eb_sel = pin;
1033 CASE_VECTOR_KSTEP_SEL(LX, pin);
1034 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
1035 NEKO_EB_SEL(LX, pin));
1036 } else {
1037 *nwf_sel = pin;
1038 CASE_VECTOR_MFMA_SEL(LX, pin);
1039 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
1040 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
1041 }
1042 log_message(neko_log_buf);
1043 log_end_section();
1044 return strat;
1045 }
1046
1047 if (strat) {
1048 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
1049 log_message(neko_log_buf);
1050 }
1051
1052 /* Formulations the sweep considers, see NEKO_TUNE_FOR(). NEKO_MFMA_TUNE
1053 keeps the matrix core variant out of an unpinned sweep, but naming it
1054 explicitly still measures it */
1055 const bool try_kstep = (strat == 0 || strat == 2);
1056 const bool try_mfma = mfma && ((strat == 3) ||
1057 (strat == 0 && neko_mfma_sweep()));
1058
1059 /* With neither the blocking sweep nor a matrix core variant there is one
1060 candidate, which is what these kernels have always run */
1061 if (!sweep && !try_mfma) {
1062 CASE_VECTOR_KSTEP(LX, 0);
1063 sprintf(neko_log_buf, "Vector Ax : %d elem/block (untuned)",
1064 NEKO_EB_SEL(LX, 0));
1065 log_message(neko_log_buf);
1066 log_end_section();
1067 return 2;
1068 }
1069
1070 HIP_CHECK(hipEventCreate(&start));
1071 HIP_CHECK(hipEventCreate(&stop));
1072
1073 /* Verify every candidate against the reference formulation on a synthetic
1074 input before anything is timed, see neko_tune_verify_begin() in
1075 elem_block_tune.h. A candidate that does not reproduce it is reported
1076 and dropped from the comparison below. */
1077 neko_tune_rej_t rej2 = {{0}}, rej3 = {{0}};
1078 {
1079 neko_tune_verify_t vf;
1080 real *vouts[] = { (real *) au, (real *) av, (real *) aw };
1081 void **vins[] = { &u, &v, &w };
1082 if (neko_tune_verify_begin(&vf, vouts, 3, vins, 3,
1083 (*nelv) * LX * LX * LX, stream)) {
1084 sprintf(neko_log_buf, "Verify : rel tol %.1e, synthetic input",
1085 vf.tol);
1086 log_message(neko_log_buf);
1087 CASE_VECTOR_KSTEP_SEL(LX, 0);
1088 neko_tune_verify_ref(&vf, stream);
1089 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1090 NEKO_TUNE_VERIFY_CAND(vf, rej2, CASE_VECTOR_KSTEP_SEL, LX, c, "KSTEP");
1091 }
1092 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1093 NEKO_TUNE_VERIFY_CAND(vf, rej3, CASE_VECTOR_MFMA_SEL, LX, c, "MFMA");
1094 }
1095 neko_tune_verify_end(&vf, stream);
1096 } else {
1097 sprintf(neko_log_buf, "Verify : %s", neko_tune_verify_on() ?
1098 "skipped, no memory for it" : "off (NEKO_TUNE_VERIFY=0)");
1099 log_message(neko_log_buf);
1100 }
1101 }
1102
1103 /* Warm every variant before timing anything, see tune_compr() */
1104 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1105 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1106 CASE_VECTOR_KSTEP_SEL(LX, c);
1107 }
1108 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1109 CASE_VECTOR_MFMA_SEL(LX, c);
1110 }
1111 }
1112
1113 /* Interleaved rounds, best time per variant, see tune_compr() */
1114 for (int r = 0; r < rounds; r++) {
1115 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1116 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_SEL, LX, c, iters);
1117 }
1118 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1119 NEKO_TUNE_TIME(time3, CASE_VECTOR_MFMA_SEL, LX, c, iters);
1120 }
1121 }
1122
1123 /* Candidates the verification pass rejected drop out here, after the
1124 timing, so the sweep's own bookkeeping is untouched */
1125 NEKO_TUNE_APPLY_REJ(time2, rej2, NEKO_EB_CANDIDATES);
1126 NEKO_TUNE_APPLY_REJ(time3, rej3, NEKO_MFMA_CANDIDATES);
1127
1128 NEKO_TUNE_LOG_KSTEP(LX, time2);
1129 NEKO_TUNE_LOG_MFMA_VEC(LX, time3);
1130
1131 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1132 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
1133 *eb_sel = best2;
1134 *nwf_sel = best3;
1135
1136 /* The mfma variant joins the comparison only where it exists, its
1137 candidates are left at NEKO_TUNE_INIT otherwise */
1138 retval = (time3[best3] < time2[best2]) ? 3 : 2;
1139
1140 HIP_CHECK(hipEventDestroy(start));
1141 HIP_CHECK(hipEventDestroy(stop));
1142
1143 /* The tuner stands in for a real Ax evaluation, and the variants do not
1144 sum in the same order, so leave the output of the chosen kernel in
1145 au, av and aw */
1146 if (retval == 2) {
1147 CASE_VECTOR_KSTEP_SEL(LX, best2);
1148 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1149 NEKO_EB_SEL(LX, best2));
1150 } else {
1151 CASE_VECTOR_MFMA_SEL(LX, best3);
1152 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
1153 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
1154 NEKO_MFMA_EB(LX, best3));
1155 }
1156 log_message(neko_log_buf);
1157 log_end_section();
1158 return retval;
1159}
1160
1161/* Padded variant of tune_compr_vector(), for the orders whose kstep kernel
1162 stages with a padded stride. The matrix core variant is the same either
1163 way */
1164template < const int LX >
1165int tune_compr_vector_padded(void *au, void *av, void *aw,
1166 void *u, void *v, void *w,
1167 void *dx, void *dy, void *dz, void *h1,
1168 void *g11, void *g22, void *g33, void *g12,
1169 void *g13, void *g23, void *ind, int *nelv,
1170 int *lx, int *eb_sel, int *nwf_sel) {
1171 hipEvent_t start, stop;
1172 float time2[NEKO_EB_CANDIDATES];
1173 int best2 = 0;
1174 float time3[NEKO_MFMA_CANDIDATES];
1175 int best3 = 0;
1176 const int rounds = neko_tune_rounds();
1177 const int iters = neko_tune_iters();
1178 const int sweep = neko_eb_sweep();
1179 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
1180 elements per block sweep is off */
1181 const int eb_cand = sweep ? NEKO_EB_CANDIDATES : 1;
1182 /* Candidates of the matrix core sweep, see neko_mfma_candidates() */
1183 const int mfma_cand = neko_mfma_candidates();
1184 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
1185 const int eb_pin = neko_eb_pin();
1186 const int nwf_pin = neko_mfma_pin();
1187 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
1188 int strat = 0;
1189 /* Hardware capability, which the explicit NEKO_AUTOTUNE=MFMA pin needs */
1190 const bool mfma = mfma_lx_supported<LX>() && hip_have_mfma();
1191 int retval;
1192
1193 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1194 time2[c] = NEKO_TUNE_INIT;
1195 }
1196 for (int c = 0; c < NEKO_MFMA_CANDIDATES; c++) {
1197 time3[c] = NEKO_TUNE_INIT;
1198 }
1199
1200 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
1201
1202 char *env_value = NULL;
1203 char neko_log_buf[80];
1204
1205 env_value = getenv("NEKO_AUTOTUNE");
1206
1207 sprintf(neko_log_buf, "Autotune Ax compr vec (lx: %d)", *lx);
1208 log_section(neko_log_buf);
1209
1210 *eb_sel = 0;
1211 *nwf_sel = 0;
1212
1213 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune_compr() */
1214 if (env_value) {
1215 if ( !strcmp(env_value, "KSTEP") || !strcmp(env_value, "1D") ) {
1216 /* No 1d variant here, so either pin lands on kstep */
1217 strat = 2;
1218 } else if ( !strcmp(env_value, "MFMA") ) {
1219 if (mfma) {
1220 strat = 3;
1221 } else {
1222 sprintf(neko_log_buf, "MFMA strategy not available for this config");
1223 log_error(neko_log_buf);
1224 }
1225 } else {
1226 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1227 log_error(neko_log_buf);
1228 }
1229 }
1230
1231 /* Formulation and geometry both pinned leaves nothing to measure */
1232 const int pin = (strat == 2) ? eb_pin : (strat == 3) ? nwf_pin : -1;
1233
1234 if (pin >= 0) {
1235 if (strat == 2) {
1236 *eb_sel = pin;
1237 CASE_VECTOR_KSTEP_PADDED_SEL(LX, pin);
1238 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
1239 NEKO_EB_SEL(LX, pin));
1240 } else {
1241 *nwf_sel = pin;
1242 CASE_VECTOR_MFMA_SEL(LX, pin);
1243 sprintf(neko_log_buf, "Set by env : 3 (MFMA %s, %d wf, %d elem/blk)",
1244 mfma_tile_name(pin), NEKO_MFMA_NWF(pin), NEKO_MFMA_EB(LX, pin));
1245 }
1246 log_message(neko_log_buf);
1247 log_end_section();
1248 return strat;
1249 }
1250
1251 if (strat) {
1252 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
1253 log_message(neko_log_buf);
1254 }
1255
1256 /* Formulations the sweep considers, see NEKO_TUNE_FOR(). NEKO_MFMA_TUNE
1257 keeps the matrix core variant out of an unpinned sweep, but naming it
1258 explicitly still measures it */
1259 const bool try_kstep = (strat == 0 || strat == 2);
1260 const bool try_mfma = mfma && ((strat == 3) ||
1261 (strat == 0 && neko_mfma_sweep()));
1262
1263 /* One candidate, see tune_compr_vector() */
1264 if (!sweep && !try_mfma) {
1265 CASE_VECTOR_KSTEP_PADDED(LX, 0);
1266 sprintf(neko_log_buf, "Vector Ax : %d elem/block (untuned)",
1267 NEKO_EB_SEL(LX, 0));
1268 log_message(neko_log_buf);
1269 log_end_section();
1270 return 2;
1271 }
1272
1273 HIP_CHECK(hipEventCreate(&start));
1274 HIP_CHECK(hipEventCreate(&stop));
1275
1276 /* Verify every candidate against the reference formulation on a synthetic
1277 input before anything is timed, see neko_tune_verify_begin() in
1278 elem_block_tune.h. A candidate that does not reproduce it is reported
1279 and dropped from the comparison below. */
1280 neko_tune_rej_t rej2 = {{0}}, rej3 = {{0}};
1281 {
1282 neko_tune_verify_t vf;
1283 real *vouts[] = { (real *) au, (real *) av, (real *) aw };
1284 void **vins[] = { &u, &v, &w };
1285 if (neko_tune_verify_begin(&vf, vouts, 3, vins, 3,
1286 (*nelv) * LX * LX * LX, stream)) {
1287 sprintf(neko_log_buf, "Verify : rel tol %.1e, synthetic input",
1288 vf.tol);
1289 log_message(neko_log_buf);
1290 CASE_VECTOR_KSTEP_PADDED_SEL(LX, 0);
1291 neko_tune_verify_ref(&vf, stream);
1292 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1293 NEKO_TUNE_VERIFY_CAND(vf, rej2, CASE_VECTOR_KSTEP_PADDED_SEL, LX, c,
1294 "KSTEP");
1295 }
1296 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1297 NEKO_TUNE_VERIFY_CAND(vf, rej3, CASE_VECTOR_MFMA_SEL, LX, c, "MFMA");
1298 }
1299 neko_tune_verify_end(&vf, stream);
1300 } else {
1301 sprintf(neko_log_buf, "Verify : %s", neko_tune_verify_on() ?
1302 "skipped, no memory for it" : "off (NEKO_TUNE_VERIFY=0)");
1303 log_message(neko_log_buf);
1304 }
1305 }
1306
1307 /* Warm every variant before timing anything, see tune_compr() */
1308 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1309 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1310 CASE_VECTOR_KSTEP_PADDED_SEL(LX, c);
1311 }
1312 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1313 CASE_VECTOR_MFMA_SEL(LX, c);
1314 }
1315 }
1316
1317 /* Interleaved rounds, best time per variant, see tune_compr() */
1318 for (int r = 0; r < rounds; r++) {
1319 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1320 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_PADDED_SEL, LX, c, iters);
1321 }
1322 NEKO_TUNE_FOR(c, try_mfma, nwf_pin, mfma_cand) {
1323 NEKO_TUNE_TIME(time3, CASE_VECTOR_MFMA_SEL, LX, c, iters);
1324 }
1325 }
1326
1327 /* Candidates the verification pass rejected drop out here, after the
1328 timing, so the sweep's own bookkeeping is untouched */
1329 NEKO_TUNE_APPLY_REJ(time2, rej2, NEKO_EB_CANDIDATES);
1330 NEKO_TUNE_APPLY_REJ(time3, rej3, NEKO_MFMA_CANDIDATES);
1331
1332 NEKO_TUNE_LOG_KSTEP(LX, time2);
1333 NEKO_TUNE_LOG_MFMA_VEC(LX, time3);
1334
1335 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1336 NEKO_TUNE_BEST(time3, best3, NEKO_MFMA_CANDIDATES);
1337 *eb_sel = best2;
1338 *nwf_sel = best3;
1339
1340 /* The mfma variant joins the comparison only where it exists, its
1341 candidates are left at NEKO_TUNE_INIT otherwise */
1342 retval = (time3[best3] < time2[best2]) ? 3 : 2;
1343
1344 HIP_CHECK(hipEventDestroy(start));
1345 HIP_CHECK(hipEventDestroy(stop));
1346
1347 /* Leave the output of the chosen kernel in au, av and aw, see
1348 tune_compr_vector() */
1349 if (retval == 2) {
1350 CASE_VECTOR_KSTEP_PADDED_SEL(LX, best2);
1351 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1352 NEKO_EB_SEL(LX, best2));
1353 } else {
1354 CASE_VECTOR_MFMA_SEL(LX, best3);
1355 sprintf(neko_log_buf, "Chose : 3 (MFMA %s, %d wf, %d elem/blk)",
1356 mfma_tile_name(best3), NEKO_MFMA_NWF(best3),
1357 NEKO_MFMA_EB(LX, best3));
1358 }
1359 log_message(neko_log_buf);
1360 log_end_section();
1361 return retval;
1362}
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:1165
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:711
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:962
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
real
double real
Definition device_config.h:12
ax_helm_compr_kernel.h
HIP_CHECK
#define HIP_CHECK(err)
Definition check.h:8
elem_block_tune.h
neko_tune_verify_ref
static void neko_tune_verify_ref(neko_tune_verify_t *v, hipStream_t stream)
Definition elem_block_tune.h:423
NEKO_TUNE_VERIFY_CAND
#define NEKO_TUNE_VERIFY_CAND(VF, REJ, LAUNCH, LX, C, NAME)
Definition elem_block_tune.h:484
neko_tune_verify_end
static void neko_tune_verify_end(neko_tune_verify_t *v, hipStream_t stream)
Definition elem_block_tune.h:471
neko_tune_verify_on
static int neko_tune_verify_on()
Definition elem_block_tune.h:283
neko_tune_verify_begin
static int neko_tune_verify_begin(neko_tune_verify_t *v, real **outs, const int nout, void ***ins, const int nin, const int n, hipStream_t stream)
Definition elem_block_tune.h:381
NEKO_TUNE_APPLY_REJ
#define NEKO_TUNE_APPLY_REJ(T, REJ, N)
Definition elem_block_tune.h:498
NEKO_TUNE_LOG_KSTEP
#define NEKO_TUNE_LOG_KSTEP(LX, T2)
Definition elem_block_tune.h:215
NEKO_TUNE_LOG_MFMA_VEC
#define NEKO_TUNE_LOG_MFMA_VEC(LX, T3)
Definition mfma_kernel.h:655
NEKO_TUNE_LOG_MFMA
#define NEKO_TUNE_LOG_MFMA(LX, T3)
Definition mfma_kernel.h:633
mfma_tile_name
static const char * mfma_tile_name(const int c)
Definition mfma_kernel.h:398
NEKO_MFMA_CANDIDATES
#define NEKO_MFMA_CANDIDATES
Definition mfma_kernel.h:278
neko_mfma_pin
static int neko_mfma_pin()
Definition mfma_kernel.h:592
neko_mfma_candidates
static int neko_mfma_candidates()
Definition mfma_kernel.h:624
hip_have_mfma
static bool hip_have_mfma()
Definition mfma_kernel.h:144
NEKO_MFMA_NWF
#define NEKO_MFMA_NWF(C)
Definition mfma_kernel.h:280
NEKO_MFMA_EB
#define NEKO_MFMA_EB(LX, C)
Definition mfma_kernel.h:472
neko_mfma_sweep
static int neko_mfma_sweep()
Definition mfma_kernel.h:570
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)
neko_tune_rej_t
Definition elem_block_tune.h:372
neko_tune_verify_t
Definition elem_block_tune.h:361