Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
ax_helm_compr.cu
Go to the documentation of this file.
1/*
2 Copyright (c) 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/*
36 * Compressed geometric factors: cuda_ax_helm_compr() and
37 * cuda_ax_helm_compr_vector() are cuda_ax_helm() and cuda_ax_helm_vector()
38 * of ax_helm.cu on the kernels of ax_helm_compr_kernel.h, with the 1-based
39 * copy index of every element, ind, passed after g23. Same strategies, same
40 * candidates, same tuner and the same environment variables (NEKO_AUTOTUNE,
41 * NEKO_EB, NEKO_EB_TUNE, NEKO_CHUNKS, NEKO_DMMA_NW, NEKO_DMMA_TMA_NW,
42 * NEKO_TUNE_ROUNDS, NEKO_TUNE_ITERS), but a tuner state of its own, since
43 * reading the compressed factors changes the memory traffic the candidates
44 * are measured under; its log sections read "Ax compr". The mass term of the
45 * vector operator is cuda_ax_helm_vector_part2() of ax_helm.cu, which reads
46 * no geometric factors.
47 *
48 * Derived mechanically from ax_helm.cu; a change to one belongs in the other.
49 */
50
51#include <string.h>
52#include <stdlib.h>
53#include <stdio.h>
54#include "ax_helm_compr_kernel.h"
55#include "elem_block_tune.h"
56#include <device/device_config.h>
57#include <device/cuda/check.h>
58
59extern "C" {
60 #include <common/neko_log.h>
61}
62
63template < const int>
64int tune_compr(void *w, void *u, void *dx, void *dy, void *dz, void *dxt,
65 void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33,
66 void *g12, void *g13, void *g23, void *ind, int *nelv, int *lx,
67 int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel);
68
69template < const int>
70int tune_compr_padded(void *w, void *u, void *dx, void *dy, void *dz, void *dxt,
71 void *dyt, void *dzt, void *h1, void *g11, void *g22,
72 void *g33, void *g12, void *g13, void *g23, void *ind,
73 int *nelv, int *lx, int *eb_sel, int *ch_sel, int *nw_sel,
74 int *tw_sel);
75
76template < const int>
77int tune_compr_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
78 void *dx, void *dy, void *dz, void *h1, void *g11,
79 void *g22, void *g33, void *g12, void *g13, void *g23,
80 void *ind, int *nelv, int *lx, int *eb_sel, int *nw_sel,
81 int *tw_sel, int *bw_sel);
82
83template < const int>
84int tune_compr_vector_padded(void *au, void *av, void *aw, void *u, void *v,
85 void *w, void *dx, void *dy, void *dz, void *h1,
86 void *g11, void *g22, void *g33, void *g12,
87 void *g13, void *g23, void *ind, int *nelv,
88 int *lx, int *eb_sel, int *nw_sel, int *tw_sel,
89 int *bw_sel);
90
91extern "C" {
92
97 void cuda_ax_helm_compr(void *w, void *u, void *dx, void *dy, void *dz,
98 void *dxt, void *dyt, void *dzt, void *h1, void *g11,
99 void *g22, void *g33, void *g12, void *g13, void *g23,
100 void *ind, int *nelv, int *lx) {
101
102 static int autotune[17] = { 0 };
103 /* Elements per block candidate chosen by the autotuner, see
104 elem_block<> in elem_block.h */
105 static int autotune_eb[17] = { 0 };
106 /* chunk candidate chosen for the 1d variant */
107 static int autotune_ch[17] = { 0 };
108 /* warps per block candidate chosen for the dmma variant */
109 static int autotune_nw[17] = { 0 };
110 /* warps per block candidate chosen for the tma staged dmma variant */
111 static int autotune_tw[17] = { 0 };
112
113 const dim3 nblcks_1d((*nelv), 1, 1);
114 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
115
116#define CASE_1D(LX, C) \
117 ax_helm_compr_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
118 <<<nblcks_1d, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
119 ((real *) w, (real *) u, \
120 (real *) dx, (real *) dy, (real *) dz, \
121 (real *) dxt, (real *) dyt, (real *) dzt, (real *) h1,\
122 (real *) g11, (real *) g22, (real *) g33, \
123 (real *) g12, (real *) g13, (real *) g23, \
124 (const int *) ind); \
125 CUDA_CHECK(cudaGetLastError());
126
127/* Runtime dispatch onto the tuned chunk candidate */
128#define CASE_1D_SEL(LX, SEL) \
129 switch (SEL) { \
130 case 0: CASE_1D(LX, 0); break; \
131 case 1: CASE_1D(LX, 1); break; \
132 case 2: CASE_1D(LX, 2); break; \
133 default: CASE_1D(LX, 3); break; \
134 }
135
136#define CASE_KSTEP(LX, C) \
137 ax_helm_compr_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
138 <<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
139 ((real *) w, (real *) u, \
140 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
141 (real *) g11, (real *) g22, (real *) g33, \
142 (real *) g12, (real *) g13, (real *) g23, \
143 (const int *) ind, *nelv); \
144 CUDA_CHECK(cudaGetLastError());
145
146#define CASE_KSTEP_PADDED(LX, C) \
147 ax_helm_compr_kernel_kstep_padded<real, LX, NEKO_EB(LX, C)> \
148 <<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
149 ((real *) w, (real *) u, \
150 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
151 (real *) g11, (real *) g22, (real *) g33, \
152 (real *) g12, (real *) g13, (real *) g23, \
153 (const int *) ind, *nelv); \
154 CUDA_CHECK(cudaGetLastError());
155
156/* Runtime dispatch onto the tuned elements per block candidate */
157#define CASE_KSTEP_SEL(LX, SEL) \
158 switch (SEL) { \
159 case 0: CASE_KSTEP(LX, 0); break; \
160 case 1: CASE_KSTEP(LX, 1); break; \
161 default: CASE_KSTEP(LX, 2); break; \
162 }
163
164#define CASE_KSTEP_PADDED_SEL(LX, SEL) \
165 switch (SEL) { \
166 case 0: CASE_KSTEP_PADDED(LX, 0); break; \
167 case 1: CASE_KSTEP_PADDED(LX, 1); break; \
168 default: CASE_KSTEP_PADDED(LX, 2); break; \
169 }
170
171#define CASE_DMMA(LX, C) \
172 ax_helm_compr_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
173 <<<NEKO_DMMA_NBLCKS(*nelv, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
174 ((real *) w, (real *) u, \
175 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
176 (real *) g11, (real *) g22, (real *) g33, \
177 (real *) g12, (real *) g13, (real *) g23, \
178 (const int *) ind, *nelv); \
179 CUDA_CHECK(cudaGetLastError());
180
181/* Runtime dispatch onto the tuned warps per block candidate */
182#define CASE_DMMA_SEL(LX, SEL) \
183 switch (SEL) { \
184 case 0: CASE_DMMA(LX, 0); break; \
185 case 1: CASE_DMMA(LX, 1); break; \
186 default: CASE_DMMA(LX, 2); break; \
187 }
188
189/*
190 * The TMA staged dmma variant. Same grid and the same warps per block
191 * candidates as CASE_DMMA -- at the one lx it supports NEKO_DMMA_PACK is 1,
192 * so the grid is nelv blocks -- but it takes no nelv: there is no packed tail
193 * to clamp when a block is exactly an element. See dmma_tma_kernel.h.
194 */
195#define CASE_DMMA_TMA(LX, C) \
196 ax_helm_compr_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
197 <<<NEKO_DMMA_NBLCKS(*nelv, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
198 ((real *) w, (real *) u, \
199 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
200 (real *) g11, (real *) g22, (real *) g33, \
201 (real *) g12, (real *) g13, (real *) g23, \
202 (const int *) ind); \
203 CUDA_CHECK(cudaGetLastError());
204
205/* Runtime dispatch onto the tuned warps per block candidate */
206#define CASE_DMMA_TMA_SEL(LX, SEL) \
207 switch (SEL) { \
208 case 0: CASE_DMMA_TMA(LX, 0); break; \
209 case 1: CASE_DMMA_TMA(LX, 1); break; \
210 default: CASE_DMMA_TMA(LX, 2); break; \
211 }
212
213#define CASE(LX) \
214 case LX: \
215 if(autotune[LX] == 0 ) { \
216 autotune[LX]=tune_compr<LX>( w, u, \
217 dx, dy, dz, \
218 dxt, dyt, dzt,h1, \
219 g11, g22, g33, \
220 g12, g13, g23, ind, nelv, lx, \
221 &autotune_eb[LX], &autotune_ch[LX], \
222 &autotune_nw[LX], &autotune_tw[LX]); \
223 } else if (autotune[LX] == 1 ) { \
224 CASE_1D_SEL(LX, autotune_ch[LX]); \
225 } else if (autotune[LX] == 2 ) { \
226 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
227 } else if (autotune[LX] == 3 ) { \
228 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
229 } else if (autotune[LX] == 4 ) { \
230 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
231 } \
232 break
233
234#define CASE_PADDED(LX) \
235 case LX: \
236 if(autotune[LX] == 0 ) { \
237 autotune[LX]=tune_compr_padded<LX>(w, u, \
238 dx, dy, dz, \
239 dxt, dyt, dzt,h1, \
240 g11, g22, g33, \
241 g12, g13, g23, ind,nelv,lx, \
242 &autotune_eb[LX], &autotune_ch[LX], \
243 &autotune_nw[LX], &autotune_tw[LX]); \
244 } else if (autotune[LX] == 1 ) { \
245 CASE_1D_SEL(LX, autotune_ch[LX]); \
246 } else if (autotune[LX] == 2 ) { \
247 CASE_KSTEP_PADDED_SEL(LX, autotune_eb[LX]); \
248 } else if (autotune[LX] == 3 ) { \
249 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
250 } else if (autotune[LX] == 4 ) { \
251 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
252 } \
253 break
254
255/*
256 * High order cases have no 1d variant to compare against (its shared memory
257 * footprint grows as LX^3), so they are not tuned and keep one element per
258 * block, i.e. candidate 0
259 */
260#define CASE_LARGE(LX) \
261 case LX: \
262 CASE_KSTEP(LX, 0); \
263 break
264
265#define CASE_LARGE_PADDED(LX) \
266 case LX: \
267 CASE_KSTEP_PADDED(LX, 0); \
268 break
269
270 if ((*lx) < 12) {
271 switch(*lx) {
272 CASE(2);
273 CASE(3);
274 CASE_PADDED(4);
275 CASE(5);
276 CASE(6);
277 CASE(7);
278 CASE_PADDED(8);
279 CASE(9);
280 CASE(10);
281 CASE(11);
282 default:
283 {
284 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
285 exit(1);
286 }
287 }
288 }
289 else {
290 switch(*lx) {
291 CASE_LARGE(12);
292 CASE_LARGE(13);
293 CASE_LARGE(14);
294 CASE_LARGE(15);
295 CASE_LARGE_PADDED(16);
296 default:
297 {
298 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
299 exit(1);
300 }
301 }
302 }
303 }
304
309 void cuda_ax_helm_compr_vector(void *au, void *av, void *aw, void *u, void *v,
310 void *w, void *dx, void *dy, void *dz,
311 void *dxt, void *dyt, void *dzt, void *h1,
312 void *g11, void *g22, void *g33, void *g12,
313 void *g13, void *g23, void *ind, int *nelv,
314 int *lx) {
315
316 /* Strategy chosen by the autotuner: 2 kstep, 3 dmma */
317 static int autotune_vec[17] = { 0 };
318 /* elements per block candidate chosen for the kstep variant */
319 static int autotune_vec_eb[17] = { 0 };
320 /* warps per block candidate chosen for the dmma variant */
321 static int autotune_vec_nw[17] = { 0 };
322 /* warps per block candidate chosen for the tma staged dmma variant */
323 static int autotune_vec_tw[17] = { 0 };
324 /* warps per block candidate chosen for the batched tma variant */
325 static int autotune_vec_bw[17] = { 0 };
326
327 const dim3 nblcks_dmma((*nelv), 1, 1);
328 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
329
330/*
331 * The vector kstep variant sweeps the same elements per block candidates as
332 * the scalar one, and the tuner picks between it and the dmma variant. The
333 * vector kernels hold three times the register blocked state -- six T[LX]
334 * arrays rather than two -- and measure at 254-255 registers on sm_90 for
335 * every lx from 8 up, spilling at 10, 12, 14 and 16, so blocking is not
336 * expected to buy much here: it does not change registers per thread, only
337 * threads per block, so it saves the derivative matrix loads and nothing
338 * else. It is swept rather than assumed because that is a measurement, and
339 * because elem_block<>'s thread clamp already keeps every candidate inside
340 * the shared memory budget at every lx.
341 */
342#define CASE_VECTOR_KSTEP(LX, C) \
343 ax_helm_compr_kernel_vector_kstep<real, LX, NEKO_EB(LX, C)> \
344 <<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
345 ((real *) au, (real *) av, (real *) aw, \
346 (real *) u, (real *) v, (real *) w, \
347 (real *) dx, (real *) dy, (real *) dz, \
348 (real *) h1, (real *) g11, (real *) g22, \
349 (real *) g33, (real *) g12, (real *) g13, \
350 (real *) g23, \
351 (const int *) ind, *nelv); \
352 CUDA_CHECK(cudaGetLastError());
353
354#define CASE_VECTOR_KSTEP_PADDED(LX, C) \
355 ax_helm_compr_kernel_vector_kstep_padded<real, LX, NEKO_EB(LX, C)> \
356 <<<NEKO_EB_NBLCKS(*nelv, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
357 ((real *) au, (real *) av, (real *) aw, \
358 (real *) u, (real *) v, (real *) w, \
359 (real *) dx, (real *) dy, (real *) dz, \
360 (real *) h1, (real *) g11, (real *) g22, \
361 (real *) g33, (real *) g12, (real *) g13, \
362 (real *) g23, \
363 (const int *) ind, *nelv); \
364 CUDA_CHECK(cudaGetLastError());
365
366/* Runtime dispatch onto the tuned elements per block candidate */
367#define CASE_VECTOR_KSTEP_SEL(LX, SEL) \
368 switch (SEL) { \
369 case 0: CASE_VECTOR_KSTEP(LX, 0); break; \
370 case 1: CASE_VECTOR_KSTEP(LX, 1); break; \
371 default: CASE_VECTOR_KSTEP(LX, 2); break; \
372 }
373
374#define CASE_VECTOR_KSTEP_PADDED_SEL(LX, SEL) \
375 switch (SEL) { \
376 case 0: CASE_VECTOR_KSTEP_PADDED(LX, 0); break; \
377 case 1: CASE_VECTOR_KSTEP_PADDED(LX, 1); break; \
378 default: CASE_VECTOR_KSTEP_PADDED(LX, 2); break; \
379 }
380
381#define CASE_VECTOR_DMMA(LX, C) \
382 ax_helm_compr_kernel_dmma_vector<real, LX, NEKO_DMMA_NW(C)> \
383 <<<nblcks_dmma, NEKO_DMMA_NTHRDS(C), 0, stream>>> \
384 ((real *) au, (real *) av, (real *) aw, \
385 (real *) u, (real *) v, (real *) w, \
386 (real *) dx, (real *) dy, (real *) dz, \
387 (real *) h1, (real *) g11, (real *) g22, \
388 (real *) g33, (real *) g12, (real *) g13, \
389 (real *) g23, \
390 (const int *) ind); \
391 CUDA_CHECK(cudaGetLastError());
392
393/* Runtime dispatch onto the tuned warps per block candidate */
394#define CASE_VECTOR_DMMA_SEL(LX, SEL) \
395 switch (SEL) { \
396 case 0: CASE_VECTOR_DMMA(LX, 0); break; \
397 case 1: CASE_VECTOR_DMMA(LX, 1); break; \
398 default: CASE_VECTOR_DMMA(LX, 2); break; \
399 }
400
401/* The TMA staged vector variant, same grid and candidates as
402 CASE_VECTOR_DMMA; see dmma_tma_kernel.h */
403#define CASE_VECTOR_DMMA_TMA(LX, C) \
404 ax_helm_compr_kernel_dmma_tma_vector<real, LX, NEKO_DMMA_NW(C)> \
405 <<<nblcks_dmma, NEKO_DMMA_NTHRDS(C), 0, stream>>> \
406 ((real *) au, (real *) av, (real *) aw, \
407 (real *) u, (real *) v, (real *) w, \
408 (real *) dx, (real *) dy, (real *) dz, \
409 (real *) h1, (real *) g11, (real *) g22, \
410 (real *) g33, (real *) g12, (real *) g13, \
411 (real *) g23, \
412 (const int *) ind); \
413 CUDA_CHECK(cudaGetLastError());
414
415/* Runtime dispatch onto the tuned warps per block candidate */
416#define CASE_VECTOR_DMMA_TMA_SEL(LX, SEL) \
417 switch (SEL) { \
418 case 0: CASE_VECTOR_DMMA_TMA(LX, 0); break; \
419 case 1: CASE_VECTOR_DMMA_TMA(LX, 1); break; \
420 default: CASE_VECTOR_DMMA_TMA(LX, 2); break; \
421 }
422
423/*
424 * The batched TMA variant. Unlike every other launch here it carries a dynamic
425 * shared memory allocation -- 54800 B, past the 48 kB a block gets for free --
426 * which the kernel has to opt into once before its first launch; see
427 * ax_helm_compr_dmma_tma_batch_optin(), which is a predictable branch after
428 * that.
429 */
430#define CASE_VECTOR_DMMA_TMA_BATCH(LX, C) \
431 (void) ax_helm_compr_dmma_tma_batch_optin<real, LX, NEKO_DMMA_NW(C)>(); \
432 ax_helm_compr_kernel_dmma_tma_batch<real, LX, NEKO_DMMA_NW(C)> \
433 <<<nblcks_dmma, NEKO_DMMA_NTHRDS(C), \
434 NEKO_DMMA_TMA_BATCH_SMEM, stream>>> \
435 ((real *) au, (real *) av, (real *) aw, \
436 (real *) u, (real *) v, (real *) w, \
437 (real *) dx, (real *) dy, (real *) dz, \
438 (real *) h1, (real *) g11, (real *) g22, \
439 (real *) g33, (real *) g12, (real *) g13, \
440 (real *) g23, \
441 (const int *) ind); \
442 CUDA_CHECK(cudaGetLastError());
443
444/* Runtime dispatch onto the tuned warps per block candidate */
445#define CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, SEL) \
446 switch (SEL) { \
447 case 0: CASE_VECTOR_DMMA_TMA_BATCH(LX, 0); break; \
448 case 1: CASE_VECTOR_DMMA_TMA_BATCH(LX, 1); break; \
449 default: CASE_VECTOR_DMMA_TMA_BATCH(LX, 2); break; \
450 }
451
452#define CASE_VECTOR(LX) \
453 case LX: \
454 if (autotune_vec[LX] == 0 ) { \
455 autotune_vec[LX] = tune_compr_vector<LX>(au, av, aw, u, v, w, dx, dy, \
456 dz, h1, g11, g22, g33, g12, \
457 g13, g23, ind, nelv, lx, \
458 &autotune_vec_eb[LX], \
459 &autotune_vec_nw[LX], \
460 &autotune_vec_tw[LX], \
461 &autotune_vec_bw[LX]); \
462 } else if (autotune_vec[LX] == 2 ) { \
463 CASE_VECTOR_KSTEP_SEL(LX, autotune_vec_eb[LX]); \
464 } else if (autotune_vec[LX] == 3 ) { \
465 CASE_VECTOR_DMMA_SEL(LX, autotune_vec_nw[LX]); \
466 } else if (autotune_vec[LX] == 4 ) { \
467 CASE_VECTOR_DMMA_TMA_SEL(LX, autotune_vec_tw[LX]); \
468 } else if (autotune_vec[LX] == 5 ) { \
469 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, autotune_vec_bw[LX]); \
470 } \
471 break
472
473#define CASE_VECTOR_PADDED(LX) \
474 case LX: \
475 if (autotune_vec[LX] == 0 ) { \
476 autotune_vec[LX] = tune_compr_vector_padded<LX>(au, av, aw, u, v, w, \
477 dx, dy, dz, h1, \
478 g11, g22, g33, \
479 g12, g13, g23, ind, \
480 nelv, lx, \
481 &autotune_vec_eb[LX], \
482 &autotune_vec_nw[LX], \
483 &autotune_vec_tw[LX], \
484 &autotune_vec_bw[LX]); \
485 } else if (autotune_vec[LX] == 2 ) { \
486 CASE_VECTOR_KSTEP_PADDED_SEL(LX, autotune_vec_eb[LX]); \
487 } else if (autotune_vec[LX] == 3 ) { \
488 CASE_VECTOR_DMMA_SEL(LX, autotune_vec_nw[LX]); \
489 } else if (autotune_vec[LX] == 4 ) { \
490 CASE_VECTOR_DMMA_TMA_SEL(LX, autotune_vec_tw[LX]); \
491 } else if (autotune_vec[LX] == 5 ) { \
492 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, autotune_vec_bw[LX]); \
493 } \
494 break
495
496 switch(*lx) {
497 CASE_VECTOR(2);
498 CASE_VECTOR(3);
499 CASE_VECTOR_PADDED(4);
500 CASE_VECTOR(5);
501 CASE_VECTOR(6);
502 CASE_VECTOR(7);
503 CASE_VECTOR_PADDED(8);
504 CASE_VECTOR(9);
505 CASE_VECTOR(10);
506 CASE_VECTOR(11);
507 CASE_VECTOR(12);
508 CASE_VECTOR(13);
509 CASE_VECTOR(14);
510 CASE_VECTOR(15);
511 CASE_VECTOR_PADDED(16);
512 default:
513 {
514 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
515 exit(1);
516 }
517 }
518 }
519}
520
521template < const int LX >
522int tune_compr(void *w, void *u, void *dx, void *dy, void *dz, void *dxt,
523 void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33,
524 void *g12, void *g13, void *g23, void *ind, int *nelv, int *lx,
525 int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel) {
526 cudaEvent_t start,stop;
527 float time1[NEKO_CHUNKS_CANDIDATES];
528 int best1 = 0;
529 float time2[NEKO_EB_CANDIDATES];
530 float time3[NEKO_DMMA_CANDIDATES];
531 int best3 = 0;
532 float time4[NEKO_DMMA_CANDIDATES];
533 int best4 = 0;
534 const int rounds = neko_tune_rounds();
535 const int iters = neko_tune_iters();
536 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
537 elements per block sweep is off */
538 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
539 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
540 const int ch_pin = neko_chunks_pin();
541 const int eb_pin = neko_eb_pin();
542 const int nw_pin = neko_dmma_pin();
543 const int tw_pin = neko_dmma_tma_pin();
544 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier tune_compr()
545 returns */
546 int strat = 0;
547 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
548 /* Whether the pointers are bulk copy aligned is a property of this call
549 rather than of the kernel, so it is checked rather than assumed, see
550 dmma_tma_aligned() in dmma_tma_kernel.h */
551 const bool tma = dmma_tma_lx_supported<LX>() && cuda_have_tma() &&
552 dmma_tma_aligned(w, u, h1, g11, g22, g33, g12, g13, g23);
553 int best = 0;
554 int retval;
555
556 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
557 time2[c] = NEKO_TUNE_INIT;
558 }
559 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
560 time1[c] = NEKO_TUNE_INIT;
561 }
562 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
563 time3[c] = NEKO_TUNE_INIT;
564 time4[c] = NEKO_TUNE_INIT;
565 }
566
567 const dim3 nblcks_1d((*nelv), 1, 1);
568 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
569
570 char *env_value = NULL;
571 char neko_log_buf[80];
572
573 env_value=getenv("NEKO_AUTOTUNE");
574
575 sprintf(neko_log_buf, "Autotune Ax compr (lx: %d)", *lx);
576 log_section(neko_log_buf);
577
578 *eb_sel = 0;
579 *ch_sel = 0;
580 *nw_sel = 0;
581 *tw_sel = 0;
582
583 /*
584 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
585 * below is narrowed to that one kernel family, but its geometry -- the
586 * chunk size, the elements per block, the warps per block -- is still
587 * measured candidate against candidate. A formulation this build or this
588 * device does not have is reported and ignored, leaving the full sweep.
589 */
590 if(env_value) {
591 if( !strcmp(env_value,"1D") ) {
592 strat = 1;
593 } else if( !strcmp(env_value,"KSTEP") ) {
594 strat = 2;
595 } else if( !strcmp(env_value,"DMMA") ) {
596 if (dmma) {
597 strat = 3;
598 } else {
599 sprintf(neko_log_buf, "DMMA strategy not available for this config");
600 log_error(neko_log_buf);
601 }
602 } else if( !strcmp(env_value,"DMMA_TMA") ) {
603 if (tma) {
604 strat = 4;
605 } else {
606 sprintf(neko_log_buf,
607 "DMMA_TMA strategy not available for this config");
608 log_error(neko_log_buf);
609 }
610 } else {
611 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
612 log_error(neko_log_buf);
613 }
614 }
615
616 /* Geometry of the pinned formulation, if its own variable fixes that too.
617 Both pinned leaves nothing to measure, so the kernel is launched once and
618 reported, which is what pinning has always done */
619 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
620 (strat == 3) ? nw_pin : (strat == 4) ? tw_pin : -1;
621
622 if (pin >= 0) {
623 switch (strat) {
624 case 1:
625 *ch_sel = pin;
626 CASE_1D_SEL(LX, pin);
627 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
628 NEKO_CHUNKS_SEL(LX, pin));
629 break;
630 case 2:
631 *eb_sel = pin;
632 CASE_KSTEP_SEL(LX, pin);
633 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
634 NEKO_EB_SEL(LX, pin));
635 break;
636 case 3:
637 *nw_sel = pin;
638 CASE_DMMA_SEL(LX, pin);
639 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
640 NEKO_DMMA_NW(pin));
641 break;
642 default:
643 *tw_sel = pin;
644 CASE_DMMA_TMA_SEL(LX, pin);
645 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
646 NEKO_DMMA_NW(pin));
647 break;
648 }
649 log_message(neko_log_buf);
650 log_end_section();
651 return strat;
652 }
653
654 if (strat) {
655 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
656 log_message(neko_log_buf);
657 }
658
659 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
660 const bool try_1d = (strat == 0 || strat == 1);
661 const bool try_kstep = (strat == 0 || strat == 2);
662 const bool try_dmma = dmma && (strat == 0 || strat == 3);
663 const bool try_tma = tma && (strat == 0 || strat == 4);
664
665 cudaEventCreate(&start);
666 cudaEventCreate(&stop);
667
668 /* Warm every variant before timing anything: each specialisation has to
669 be resident and the clocks at steady state, or whichever variant is
670 timed first is measured on a colder part */
671 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
672 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
673 CASE_1D_SEL(LX, c);
674 }
675 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
676 CASE_KSTEP_SEL(LX, c);
677 }
678 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
679 CASE_DMMA_SEL(LX, c);
680 }
681 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
682 CASE_DMMA_TMA_SEL(LX, c);
683 }
684 }
685
686 /* Interleaved rounds, best time per variant: timing them one after another
687 in a fixed order lets clock drift bias the comparison by position */
688 for (int r = 0; r < rounds; r++) {
689 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
690 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
691 }
692 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
693 NEKO_TUNE_TIME(time2, CASE_KSTEP_SEL, LX, c, iters);
694 }
695 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
696 NEKO_TUNE_TIME(time3, CASE_DMMA_SEL, LX, c, iters);
697 }
698 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
699 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA_SEL, LX, c, iters);
700 }
701 }
702
703 NEKO_TUNE_LOG(LX, time1, time2);
704 NEKO_TUNE_LOG_DMMA(LX, time3);
705 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
706
707 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
708 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
709 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
710 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
711 *eb_sel = best;
712 *ch_sel = best1;
713 *nw_sel = best3;
714 *tw_sel = best4;
715
716 if (time1[best1] < time2[best]) {
717 retval = 1;
718 } else {
719 retval = 2;
720 }
721
722 /* The dmma variants join the comparison only where they exist, their
723 candidates are left at NEKO_TUNE_INIT otherwise */
724 float tbest = (retval == 1) ? time1[best1] : time2[best];
725
726 if (time3[best3] < tbest) {
727 retval = 3;
728 tbest = time3[best3];
729 }
730 if (time4[best4] < tbest) {
731 retval = 4;
732 }
733
734 cudaEventDestroy(start);
735 cudaEventDestroy(stop);
736
737 /* The tuner stands in for a real Ax evaluation, and the variants do not
738 sum in the same order, so leave the output of the chosen kernel in w */
739 if (retval == 1) {
740 CASE_1D_SEL(LX, best1);
741 } else if (retval == 2) {
742 CASE_KSTEP_SEL(LX, best);
743 } else if (retval == 3) {
744 CASE_DMMA_SEL(LX, best3);
745 } else {
746 CASE_DMMA_TMA_SEL(LX, best4);
747 }
748
749 if (retval == 1) {
750 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
751 NEKO_CHUNKS_SEL(LX, best1));
752 } else if (retval == 2) {
753 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
754 NEKO_EB_SEL(LX, best));
755 } else if (retval == 3) {
756 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
757 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
758 } else {
759 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
760 NEKO_DMMA_NW(best4));
761 }
762 log_message(neko_log_buf);
763 log_end_section();
764 return retval;
765}
766
767template < const int LX >
768int tune_compr_padded(void *w, void *u, void *dx, void *dy, void *dz, void *dxt,
769 void *dyt, void *dzt, void *h1, void *g11, void *g22,
770 void *g33, void *g12, void *g13, void *g23, void *ind,
771 int *nelv, int *lx, int *eb_sel, int *ch_sel, int *nw_sel,
772 int *tw_sel) {
773 cudaEvent_t start, stop;
774 float time1[NEKO_CHUNKS_CANDIDATES];
775 int best1 = 0;
776 float time2[NEKO_EB_CANDIDATES];
777 float time3[NEKO_DMMA_CANDIDATES];
778 int best3 = 0;
779 float time4[NEKO_DMMA_CANDIDATES];
780 int best4 = 0;
781 const int rounds = neko_tune_rounds();
782 const int iters = neko_tune_iters();
783 /* Candidates of the kstep sweep, see tune_compr() */
784 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
785 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
786 const int ch_pin = neko_chunks_pin();
787 const int eb_pin = neko_eb_pin();
788 const int nw_pin = neko_dmma_pin();
789 const int tw_pin = neko_dmma_tma_pin();
790 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
791 int strat = 0;
792 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
793 /* Whether the pointers are bulk copy aligned is a property of this call
794 rather than of the kernel, so it is checked rather than assumed, see
795 dmma_tma_aligned() in dmma_tma_kernel.h */
796 const bool tma = dmma_tma_lx_supported<LX>() && cuda_have_tma() &&
797 dmma_tma_aligned(w, u, h1, g11, g22, g33, g12, g13, g23);
798 int best = 0;
799 int retval;
800
801 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
802 time2[c] = NEKO_TUNE_INIT;
803 }
804 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
805 time1[c] = NEKO_TUNE_INIT;
806 }
807 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
808 time3[c] = NEKO_TUNE_INIT;
809 time4[c] = NEKO_TUNE_INIT;
810 }
811
812 const dim3 nblcks_1d((*nelv), 1, 1);
813 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
814
815 char *env_value = NULL;
816 char neko_log_buf[80];
817
818 env_value=getenv("NEKO_AUTOTUNE");
819
820 sprintf(neko_log_buf, "Autotune Ax compr (lx: %d)", *lx);
821 log_section(neko_log_buf);
822
823 *eb_sel = 0;
824 *ch_sel = 0;
825 *nw_sel = 0;
826 *tw_sel = 0;
827
828 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune_compr() */
829 if(env_value) {
830 if( !strcmp(env_value,"1D") ) {
831 strat = 1;
832 } else if( !strcmp(env_value,"KSTEP") ) {
833 strat = 2;
834 } else if( !strcmp(env_value,"DMMA") ) {
835 if (dmma) {
836 strat = 3;
837 } else {
838 sprintf(neko_log_buf, "DMMA strategy not available for this config");
839 log_error(neko_log_buf);
840 }
841 } else if( !strcmp(env_value,"DMMA_TMA") ) {
842 if (tma) {
843 strat = 4;
844 } else {
845 sprintf(neko_log_buf,
846 "DMMA_TMA strategy not available for this config");
847 log_error(neko_log_buf);
848 }
849 } else {
850 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
851 log_error(neko_log_buf);
852 }
853 }
854
855 /* Formulation and geometry both pinned leaves nothing to measure */
856 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
857 (strat == 3) ? nw_pin : (strat == 4) ? tw_pin : -1;
858
859 if (pin >= 0) {
860 switch (strat) {
861 case 1:
862 *ch_sel = pin;
863 CASE_1D_SEL(LX, pin);
864 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
865 NEKO_CHUNKS_SEL(LX, pin));
866 break;
867 case 2:
868 *eb_sel = pin;
869 CASE_KSTEP_PADDED_SEL(LX, pin);
870 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
871 NEKO_EB_SEL(LX, pin));
872 break;
873 case 3:
874 *nw_sel = pin;
875 CASE_DMMA_SEL(LX, pin);
876 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
877 NEKO_DMMA_NW(pin));
878 break;
879 default:
880 *tw_sel = pin;
881 CASE_DMMA_TMA_SEL(LX, pin);
882 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
883 NEKO_DMMA_NW(pin));
884 break;
885 }
886 log_message(neko_log_buf);
887 log_end_section();
888 return strat;
889 }
890
891 if (strat) {
892 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
893 log_message(neko_log_buf);
894 }
895
896 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
897 const bool try_1d = (strat == 0 || strat == 1);
898 const bool try_kstep = (strat == 0 || strat == 2);
899 const bool try_dmma = dmma && (strat == 0 || strat == 3);
900 const bool try_tma = tma && (strat == 0 || strat == 4);
901
902 cudaEventCreate(&start);
903 cudaEventCreate(&stop);
904
905 /* Warm every variant before timing anything: each specialisation has to
906 be resident and the clocks at steady state, or whichever variant is
907 timed first is measured on a colder part */
908 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
909 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
910 CASE_1D_SEL(LX, c);
911 }
912 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
913 CASE_KSTEP_PADDED_SEL(LX, c);
914 }
915 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
916 CASE_DMMA_SEL(LX, c);
917 }
918 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
919 CASE_DMMA_TMA_SEL(LX, c);
920 }
921 }
922
923 /* Interleaved rounds, best time per variant: timing them one after another
924 in a fixed order lets clock drift bias the comparison by position */
925 for (int r = 0; r < rounds; r++) {
926 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
927 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
928 }
929 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
930 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED_SEL, LX, c, iters);
931 }
932 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
933 NEKO_TUNE_TIME(time3, CASE_DMMA_SEL, LX, c, iters);
934 }
935 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
936 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA_SEL, LX, c, iters);
937 }
938 }
939
940 NEKO_TUNE_LOG(LX, time1, time2);
941 NEKO_TUNE_LOG_DMMA(LX, time3);
942 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
943
944 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
945 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
946 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
947 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
948 *eb_sel = best;
949 *ch_sel = best1;
950 *nw_sel = best3;
951 *tw_sel = best4;
952
953 if (time1[best1] < time2[best]) {
954 retval = 1;
955 } else {
956 retval = 2;
957 }
958
959 /* The dmma variants join the comparison only where they exist, their
960 candidates are left at NEKO_TUNE_INIT otherwise */
961 float tbest = (retval == 1) ? time1[best1] : time2[best];
962
963 if (time3[best3] < tbest) {
964 retval = 3;
965 tbest = time3[best3];
966 }
967 if (time4[best4] < tbest) {
968 retval = 4;
969 }
970
971 cudaEventDestroy(start);
972 cudaEventDestroy(stop);
973
974 /* Leave the output of the chosen kernel in w, see tune_compr() */
975 if (retval == 1) {
976 CASE_1D_SEL(LX, best1);
977 } else if (retval == 2) {
978 CASE_KSTEP_PADDED_SEL(LX, best);
979 } else if (retval == 3) {
980 CASE_DMMA_SEL(LX, best3);
981 } else {
982 CASE_DMMA_TMA_SEL(LX, best4);
983 }
984
985 if (retval == 1) {
986 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
987 NEKO_CHUNKS_SEL(LX, best1));
988 } else if (retval == 2) {
989 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
990 NEKO_EB_SEL(LX, best));
991 } else if (retval == 3) {
992 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
993 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
994 } else {
995 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
996 NEKO_DMMA_NW(best4));
997 }
998 log_message(neko_log_buf);
999 log_end_section();
1000 return retval;
1001}
1002
1003
1004/*
1005 * Autotuner for the vector Ax
1006 *
1007 * The vector operator has no 1d variant, so the sweep is the elements per
1008 * block candidates of the kstep variant against the warps per block
1009 * candidates of the dmma variant -- the same two dimensions the scalar tuner
1010 * sweeps, and the same shape as the HIP vector tuner. See the note at the top
1011 * of ax_helm_compr_kernel.h for why blocking is not expected to pay here, and
1012 * why it is measured rather than assumed. Sampling follows the scalar tuner:
1013 * interleaved rounds, best round per candidate.
1014 *
1015 * @note The section title has to fit the 30 character header field in
1016 * log_section(), which truncates rather than wraps -- hence "Ax compr vec"
1017 * rather than the "Ax helm compr vector" it reads as.
1018 */
1019template < const int LX >
1020int tune_compr_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
1021 void *dx, void *dy, void *dz, void *h1, void *g11,
1022 void *g22, void *g33, void *g12, void *g13, void *g23,
1023 void *ind, int *nelv, int *lx, int *eb_sel, int *nw_sel,
1024 int *tw_sel, int *bw_sel) {
1025 cudaEvent_t start, stop;
1026 float time2[NEKO_EB_CANDIDATES];
1027 int best2 = 0;
1028 float time3[NEKO_DMMA_CANDIDATES];
1029 int best3 = 0;
1030 float time4[NEKO_DMMA_CANDIDATES];
1031 int best4 = 0;
1032 float time5[NEKO_DMMA_CANDIDATES];
1033 int best5 = 0;
1034 const int rounds = neko_tune_rounds();
1035 const int iters = neko_tune_iters();
1036 /* Candidates of the kstep sweep, see tune_compr() */
1037 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
1038 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
1039 const int eb_pin = neko_eb_pin();
1040 const int nw_pin = neko_dmma_pin();
1041 const int tw_pin = neko_dmma_tma_pin();
1042 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
1043 int strat = 0;
1044 const bool dmma = dmma_vector_lx_supported<LX>() && cuda_have_dmma();
1045 /* Whether the pointers are bulk copy aligned is a property of this call
1046 rather than of the kernel, see dmma_tma_vector_aligned() */
1047 const bool tma = dmma_tma_vector_lx_supported<LX>() && cuda_have_tma() &&
1048 dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
1049 g11, g22, g33, g12, g13, g23);
1050 /* Same pointers, and additionally the device has to grant the block the
1051 dynamic allocation, see cuda_have_tma_batch() */
1052 const bool batch = dmma_tma_batch_lx_supported<LX>() && cuda_have_tma_batch()
1053 && dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
1054 g11, g22, g33,
1055 g12, g13, g23);
1056 int retval;
1057
1058 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1059 time2[c] = NEKO_TUNE_INIT;
1060 }
1061 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
1062 time3[c] = NEKO_TUNE_INIT;
1063 time4[c] = NEKO_TUNE_INIT;
1064 time5[c] = NEKO_TUNE_INIT;
1065 }
1066
1067 const dim3 nblcks_dmma((*nelv), 1, 1);
1068 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
1069
1070 char *env_value = NULL;
1071 char neko_log_buf[80];
1072
1073 env_value=getenv("NEKO_AUTOTUNE");
1074
1075 sprintf(neko_log_buf, "Autotune Ax compr vec (lx: %d)", *lx);
1076 log_section(neko_log_buf);
1077
1078 *eb_sel = 0;
1079 *nw_sel = 0;
1080 *tw_sel = 0;
1081 *bw_sel = 0;
1082
1083 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune_compr() */
1084 if(env_value) {
1085 if( !strcmp(env_value,"KSTEP") || !strcmp(env_value,"1D") ) {
1086 /* No 1d variant here, so either pin lands on kstep */
1087 strat = 2;
1088 } else if( !strcmp(env_value,"DMMA") ) {
1089 if (dmma) {
1090 strat = 3;
1091 } else {
1092 sprintf(neko_log_buf, "DMMA strategy not available for this config");
1093 log_error(neko_log_buf);
1094 }
1095 } else if( !strcmp(env_value,"DMMA_TMA") ) {
1096 if (tma) {
1097 strat = 4;
1098 } else {
1099 sprintf(neko_log_buf,
1100 "DMMA_TMA strategy not available for this config");
1101 log_error(neko_log_buf);
1102 }
1103 } else if( !strcmp(env_value,"DMMA_TMA_BATCH") ) {
1104 if (batch) {
1105 strat = 5;
1106 } else {
1107 sprintf(neko_log_buf,
1108 "DMMA_TMA_BATCH strategy not available for this config");
1109 log_error(neko_log_buf);
1110 }
1111 } else {
1112 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1113 log_error(neko_log_buf);
1114 }
1115 }
1116
1117 /* Formulation and geometry both pinned leaves nothing to measure. The two
1118 TMA forms share NEKO_DMMA_TMA_NW, hence the same pin for both */
1119 const int pin = (strat == 2) ? eb_pin : (strat == 3) ? nw_pin :
1120 (strat == 4 || strat == 5) ? tw_pin : -1;
1121
1122 if (pin >= 0) {
1123 switch (strat) {
1124 case 2:
1125 *eb_sel = pin;
1126 CASE_VECTOR_KSTEP_SEL(LX, pin);
1127 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
1128 NEKO_EB_SEL(LX, pin));
1129 break;
1130 case 3:
1131 *nw_sel = pin;
1132 CASE_VECTOR_DMMA_SEL(LX, pin);
1133 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
1134 NEKO_DMMA_NW(pin));
1135 break;
1136 case 4:
1137 *tw_sel = pin;
1138 CASE_VECTOR_DMMA_TMA_SEL(LX, pin);
1139 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
1140 NEKO_DMMA_NW(pin));
1141 break;
1142 default:
1143 *bw_sel = pin;
1144 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, pin);
1145 sprintf(neko_log_buf, "Set by env : 5 (DMMA_TMA_BATCH, %d warps)",
1146 NEKO_DMMA_NW(pin));
1147 break;
1148 }
1149 log_message(neko_log_buf);
1150 log_end_section();
1151 return strat;
1152 }
1153
1154 if (strat) {
1155 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
1156 log_message(neko_log_buf);
1157 }
1158
1159 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
1160 const bool try_kstep = (strat == 0 || strat == 2);
1161 const bool try_dmma = dmma && (strat == 0 || strat == 3);
1162 const bool try_tma = tma && (strat == 0 || strat == 4);
1163 const bool try_batch = batch && (strat == 0 || strat == 5);
1164
1165 cudaEventCreate(&start);
1166 cudaEventCreate(&stop);
1167
1168 /* Warm every variant before timing anything, see tune_compr() */
1169 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1170 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1171 CASE_VECTOR_KSTEP_SEL(LX, c);
1172 }
1173 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
1174 CASE_VECTOR_DMMA_SEL(LX, c);
1175 }
1176 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
1177 CASE_VECTOR_DMMA_TMA_SEL(LX, c);
1178 }
1179 NEKO_TUNE_FOR(c, try_batch, tw_pin, NEKO_DMMA_CANDIDATES) {
1180 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, c);
1181 }
1182 }
1183
1184 for (int r = 0; r < rounds; r++) {
1185 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1186 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_SEL, LX, c, iters);
1187 }
1188 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
1189 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA_SEL, LX, c, iters);
1190 }
1191 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
1192 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA_SEL, LX, c, iters);
1193 }
1194 NEKO_TUNE_FOR(c, try_batch, tw_pin, NEKO_DMMA_CANDIDATES) {
1195 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH_SEL, LX, c, iters);
1196 }
1197 }
1198
1199 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1200 if (time2[c] >= NEKO_TUNE_INIT) { continue; }
1201 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call",
1202 NEKO_EB_SEL(LX, c), NEKO_TUNE_US(time2[c], iters));
1203 log_message(neko_log_buf);
1204 }
1205 NEKO_TUNE_LOG_DMMA(LX, time3);
1206 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
1207 NEKO_TUNE_LOG_DMMA_TMA_BATCH(LX, time5);
1208
1209 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1210 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
1211 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
1212 NEKO_TUNE_BEST(time5, best5, NEKO_DMMA_CANDIDATES);
1213 *eb_sel = best2;
1214 *nw_sel = best3;
1215 *tw_sel = best4;
1216 *bw_sel = best5;
1217
1218 /* The dmma variants join the comparison only where they exist, their
1219 candidates are left at NEKO_TUNE_INIT otherwise */
1220 float tbest = time2[best2];
1221
1222 retval = 2;
1223 if (time3[best3] < tbest) { retval = 3; tbest = time3[best3]; }
1224 if (time4[best4] < tbest) { retval = 4; tbest = time4[best4]; }
1225 if (time5[best5] < tbest) { retval = 5; }
1226
1227 cudaEventDestroy(start);
1228 cudaEventDestroy(stop);
1229
1230 /* The tuner stands in for a real Ax evaluation, and the variants do not
1231 sum in the same order, so leave the output of the chosen kernel in
1232 au, av and aw */
1233 if (retval == 2) {
1234 CASE_VECTOR_KSTEP_SEL(LX, best2);
1235 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1236 NEKO_EB_SEL(LX, best2));
1237 } else if (retval == 3) {
1238 CASE_VECTOR_DMMA_SEL(LX, best3);
1239 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
1240 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
1241 } else if (retval == 4) {
1242 CASE_VECTOR_DMMA_TMA_SEL(LX, best4);
1243 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
1244 NEKO_DMMA_NW(best4));
1245 } else {
1246 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, best5);
1247 sprintf(neko_log_buf, "Chose : 5 (DMMA_TMA_BATCH, %d warps)",
1248 NEKO_DMMA_NW(best5));
1249 }
1250 log_message(neko_log_buf);
1251 log_end_section();
1252 return retval;
1253}
1254
1255/* Padded variant of tune_compr_vector(), see the note on bank conflicts
1256 above */
1257template < const int LX >
1258int tune_compr_vector_padded(void *au, void *av, void *aw, void *u, void *v,
1259 void *w, void *dx, void *dy, void *dz, void *h1,
1260 void *g11, void *g22, void *g33, void *g12,
1261 void *g13, void *g23, void *ind, int *nelv,
1262 int *lx, int *eb_sel, int *nw_sel, int *tw_sel,
1263 int *bw_sel) {
1264 cudaEvent_t start, stop;
1265 float time2[NEKO_EB_CANDIDATES];
1266 int best2 = 0;
1267 float time3[NEKO_DMMA_CANDIDATES];
1268 int best3 = 0;
1269 float time4[NEKO_DMMA_CANDIDATES];
1270 int best4 = 0;
1271 float time5[NEKO_DMMA_CANDIDATES];
1272 int best5 = 0;
1273 const int rounds = neko_tune_rounds();
1274 const int iters = neko_tune_iters();
1275 /* Candidates of the kstep sweep, see tune_compr() */
1276 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
1277 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
1278 const int eb_pin = neko_eb_pin();
1279 const int nw_pin = neko_dmma_pin();
1280 const int tw_pin = neko_dmma_tma_pin();
1281 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
1282 int strat = 0;
1283 const bool dmma = dmma_vector_lx_supported<LX>() && cuda_have_dmma();
1284 /* Whether the pointers are bulk copy aligned is a property of this call
1285 rather than of the kernel, see dmma_tma_vector_aligned() */
1286 const bool tma = dmma_tma_vector_lx_supported<LX>() && cuda_have_tma() &&
1287 dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
1288 g11, g22, g33, g12, g13, g23);
1289 /* Same pointers, and additionally the device has to grant the block the
1290 dynamic allocation, see cuda_have_tma_batch() */
1291 const bool batch = dmma_tma_batch_lx_supported<LX>() && cuda_have_tma_batch()
1292 && dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
1293 g11, g22, g33,
1294 g12, g13, g23);
1295 int retval;
1296
1297 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1298 time2[c] = NEKO_TUNE_INIT;
1299 }
1300 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
1301 time3[c] = NEKO_TUNE_INIT;
1302 time4[c] = NEKO_TUNE_INIT;
1303 time5[c] = NEKO_TUNE_INIT;
1304 }
1305
1306 const dim3 nblcks_dmma((*nelv), 1, 1);
1307 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
1308
1309 char *env_value = NULL;
1310 char neko_log_buf[80];
1311
1312 env_value=getenv("NEKO_AUTOTUNE");
1313
1314 sprintf(neko_log_buf, "Autotune Ax compr vec (lx: %d)", *lx);
1315 log_section(neko_log_buf);
1316
1317 *eb_sel = 0;
1318 *nw_sel = 0;
1319 *tw_sel = 0;
1320 *bw_sel = 0;
1321
1322 /* NEKO_AUTOTUNE names a formulation and nothing more, see tune_compr() */
1323 if(env_value) {
1324 if( !strcmp(env_value,"KSTEP") || !strcmp(env_value,"1D") ) {
1325 /* No 1d variant here, so either pin lands on kstep */
1326 strat = 2;
1327 } else if( !strcmp(env_value,"DMMA") ) {
1328 if (dmma) {
1329 strat = 3;
1330 } else {
1331 sprintf(neko_log_buf, "DMMA strategy not available for this config");
1332 log_error(neko_log_buf);
1333 }
1334 } else if( !strcmp(env_value,"DMMA_TMA") ) {
1335 if (tma) {
1336 strat = 4;
1337 } else {
1338 sprintf(neko_log_buf,
1339 "DMMA_TMA strategy not available for this config");
1340 log_error(neko_log_buf);
1341 }
1342 } else if( !strcmp(env_value,"DMMA_TMA_BATCH") ) {
1343 if (batch) {
1344 strat = 5;
1345 } else {
1346 sprintf(neko_log_buf,
1347 "DMMA_TMA_BATCH strategy not available for this config");
1348 log_error(neko_log_buf);
1349 }
1350 } else {
1351 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1352 log_error(neko_log_buf);
1353 }
1354 }
1355
1356 /* Formulation and geometry both pinned leaves nothing to measure. The two
1357 TMA forms share NEKO_DMMA_TMA_NW, hence the same pin for both */
1358 const int pin = (strat == 2) ? eb_pin : (strat == 3) ? nw_pin :
1359 (strat == 4 || strat == 5) ? tw_pin : -1;
1360
1361 if (pin >= 0) {
1362 switch (strat) {
1363 case 2:
1364 *eb_sel = pin;
1365 CASE_VECTOR_KSTEP_PADDED_SEL(LX, pin);
1366 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
1367 NEKO_EB_SEL(LX, pin));
1368 break;
1369 case 3:
1370 *nw_sel = pin;
1371 CASE_VECTOR_DMMA_SEL(LX, pin);
1372 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
1373 NEKO_DMMA_NW(pin));
1374 break;
1375 case 4:
1376 *tw_sel = pin;
1377 CASE_VECTOR_DMMA_TMA_SEL(LX, pin);
1378 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
1379 NEKO_DMMA_NW(pin));
1380 break;
1381 default:
1382 *bw_sel = pin;
1383 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, pin);
1384 sprintf(neko_log_buf, "Set by env : 5 (DMMA_TMA_BATCH, %d warps)",
1385 NEKO_DMMA_NW(pin));
1386 break;
1387 }
1388 log_message(neko_log_buf);
1389 log_end_section();
1390 return strat;
1391 }
1392
1393 if (strat) {
1394 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
1395 log_message(neko_log_buf);
1396 }
1397
1398 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
1399 const bool try_kstep = (strat == 0 || strat == 2);
1400 const bool try_dmma = dmma && (strat == 0 || strat == 3);
1401 const bool try_tma = tma && (strat == 0 || strat == 4);
1402 const bool try_batch = batch && (strat == 0 || strat == 5);
1403
1404 cudaEventCreate(&start);
1405 cudaEventCreate(&stop);
1406
1407 /* Warm every variant before timing anything, see tune_compr() */
1408 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1409 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1410 CASE_VECTOR_KSTEP_PADDED_SEL(LX, c);
1411 }
1412 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
1413 CASE_VECTOR_DMMA_SEL(LX, c);
1414 }
1415 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
1416 CASE_VECTOR_DMMA_TMA_SEL(LX, c);
1417 }
1418 NEKO_TUNE_FOR(c, try_batch, tw_pin, NEKO_DMMA_CANDIDATES) {
1419 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, c);
1420 }
1421 }
1422
1423 for (int r = 0; r < rounds; r++) {
1424 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
1425 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_PADDED_SEL, LX, c, iters);
1426 }
1427 NEKO_TUNE_FOR(c, try_dmma, nw_pin, NEKO_DMMA_CANDIDATES) {
1428 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA_SEL, LX, c, iters);
1429 }
1430 NEKO_TUNE_FOR(c, try_tma, tw_pin, NEKO_DMMA_CANDIDATES) {
1431 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA_SEL, LX, c, iters);
1432 }
1433 NEKO_TUNE_FOR(c, try_batch, tw_pin, NEKO_DMMA_CANDIDATES) {
1434 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH_SEL, LX, c, iters);
1435 }
1436 }
1437
1438 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1439 if (time2[c] >= NEKO_TUNE_INIT) { continue; }
1440 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call",
1441 NEKO_EB_SEL(LX, c), NEKO_TUNE_US(time2[c], iters));
1442 log_message(neko_log_buf);
1443 }
1444 NEKO_TUNE_LOG_DMMA(LX, time3);
1445 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
1446 NEKO_TUNE_LOG_DMMA_TMA_BATCH(LX, time5);
1447
1448 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1449 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
1450 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
1451 NEKO_TUNE_BEST(time5, best5, NEKO_DMMA_CANDIDATES);
1452 *eb_sel = best2;
1453 *nw_sel = best3;
1454 *tw_sel = best4;
1455 *bw_sel = best5;
1456
1457 /* The dmma variants join the comparison only where they exist, their
1458 candidates are left at NEKO_TUNE_INIT otherwise */
1459 float tbest = time2[best2];
1460
1461 retval = 2;
1462 if (time3[best3] < tbest) { retval = 3; tbest = time3[best3]; }
1463 if (time4[best4] < tbest) { retval = 4; tbest = time4[best4]; }
1464 if (time5[best5] < tbest) { retval = 5; }
1465
1466 cudaEventDestroy(start);
1467 cudaEventDestroy(stop);
1468
1469 /* Leave the output of the chosen kernel in au, av and aw, see tune_compr() */
1470 if (retval == 2) {
1471 CASE_VECTOR_KSTEP_PADDED_SEL(LX, best2);
1472 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1473 NEKO_EB_SEL(LX, best2));
1474 } else if (retval == 3) {
1475 CASE_VECTOR_DMMA_SEL(LX, best3);
1476 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
1477 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
1478 } else if (retval == 4) {
1479 CASE_VECTOR_DMMA_TMA_SEL(LX, best4);
1480 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
1481 NEKO_DMMA_NW(best4));
1482 } else {
1483 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, best5);
1484 sprintf(neko_log_buf, "Chose : 5 (DMMA_TMA_BATCH, %d warps)",
1485 NEKO_DMMA_NW(best5));
1486 }
1487 log_message(neko_log_buf);
1488 log_end_section();
1489 return retval;
1490}
CASE_VECTOR
#define CASE_VECTOR(LX)
cuda_ax_helm_compr_vector
void cuda_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.cu:309
CASE_DMMA_SEL
#define CASE_DMMA_SEL(LX, SEL)
CASE_KSTEP_SEL
#define CASE_KSTEP_SEL(LX, SEL)
CASE
#define CASE(LX)
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 *nw_sel, int *tw_sel, int *bw_sel)
Definition ax_helm_compr.cu:1258
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 *nw_sel, int *tw_sel)
Definition ax_helm_compr.cu:522
CASE_1D_SEL
#define CASE_1D_SEL(LX, SEL)
CASE_PADDED
#define CASE_PADDED(LX)
CASE_KSTEP_PADDED_SEL
#define CASE_KSTEP_PADDED_SEL(LX, SEL)
CASE_LARGE
#define CASE_LARGE(LX)
CASE_DMMA_TMA_SEL
#define CASE_DMMA_TMA_SEL(LX, SEL)
CASE_LARGE_PADDED
#define CASE_LARGE_PADDED(LX)
cuda_ax_helm_compr
void cuda_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.cu:97
CASE_VECTOR_DMMA_TMA_SEL
#define CASE_VECTOR_DMMA_TMA_SEL(LX, SEL)
CASE_VECTOR_KSTEP_SEL
#define CASE_VECTOR_KSTEP_SEL(LX, SEL)
CASE_VECTOR_PADDED
#define CASE_VECTOR_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 *nw_sel, int *tw_sel)
Definition ax_helm_compr.cu:768
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 *nw_sel, int *tw_sel, int *bw_sel)
Definition ax_helm_compr.cu:1020
CASE_VECTOR_KSTEP_PADDED_SEL
#define CASE_VECTOR_KSTEP_PADDED_SEL(LX, SEL)
CASE_VECTOR_DMMA_TMA_BATCH_SEL
#define CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, SEL)
CASE_VECTOR_DMMA_SEL
#define CASE_VECTOR_DMMA_SEL(LX, SEL)
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_US
#define NEKO_TUNE_US(T, ITERS)
Definition elem_block_tune.h:183
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
NEKO_DMMA_CANDIDATES
#define NEKO_DMMA_CANDIDATES
Definition dmma_kernel.h:123
cuda_have_dmma
static bool cuda_have_dmma()
Definition dmma_kernel.h:230
NEKO_DMMA_PACK
#define NEKO_DMMA_PACK(LX)
Definition dmma_kernel.h:148
NEKO_TUNE_LOG_DMMA
#define NEKO_TUNE_LOG_DMMA(LX, T3)
Definition dmma_kernel.h:175
NEKO_DMMA_NW
#define NEKO_DMMA_NW(C)
Definition dmma_kernel.h:124
neko_dmma_pin
static int neko_dmma_pin()
Definition dmma_kernel.h:155
NEKO_TUNE_LOG_DMMA_TMA_BATCH
#define NEKO_TUNE_LOG_DMMA_TMA_BATCH(LX, T5)
Definition dmma_tma_kernel.h:572
dmma_tma_aligned
static bool dmma_tma_aligned(const void *w, const void *u, const void *h1, const void *g11, const void *g22, const void *g33, const void *g12, const void *g13, const void *g23)
Definition dmma_tma_kernel.h:455
dmma_tma_vector_aligned
static bool dmma_tma_vector_aligned(const void *au, const void *av, const void *aw, const void *u, const void *v, const void *w, const void *h1, const void *g11, const void *g22, const void *g33, const void *g12, const void *g13, const void *g23)
Definition dmma_tma_kernel.h:469
cuda_have_tma_batch
static bool cuda_have_tma_batch()
Definition dmma_tma_kernel.h:337
NEKO_TUNE_LOG_DMMA_TMA
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
Definition dmma_tma_kernel.h:585
neko_dmma_tma_pin
static int neko_dmma_tma_pin()
Definition dmma_tma_kernel.h:555
cuda_have_tma
static bool cuda_have_tma()
Definition dmma_tma_kernel.h:165
ax_helm_compr_kernel.h
elem_block_tune.h
neko_log.h
log_error
void log_error(char *msg)
log_message
void log_message(char *msg)
log_end_section
void log_end_section()
log_section
void log_section(char *msg)