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