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 const int sweep = neko_eb_sweep();
529 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
530 /* Whether the pointers are bulk copy aligned is a property of this call
531 rather than of the kernel, so it is checked rather than assumed, see
532 dmma_tma_aligned() in dmma_tma_kernel.h */
533 const bool tma = dmma_tma_lx_supported<LX>() && cuda_have_tma() &&
534 dmma_tma_aligned(w, u, h1, g11, g22, g33, g12, g13, g23);
535 int best = 0;
536 int retval;
537
538 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
539 time2[c] = NEKO_TUNE_INIT;
540 }
541 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
542 time1[c] = NEKO_TUNE_INIT;
543 }
544 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
545 time3[c] = NEKO_TUNE_INIT;
546 time4[c] = NEKO_TUNE_INIT;
547 }
548
549 const dim3 nblcks_1d((*nelv), 1, 1);
550 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
551
552 char *env_value = NULL;
553 char neko_log_buf[80];
554
555 env_value=getenv("NEKO_AUTOTUNE");
556
557 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
558 log_section(neko_log_buf);
559
560 *eb_sel = 0;
561 *ch_sel = 0;
562 *tw_sel = 0;
563
564 if(env_value) {
565 if( !strcmp(env_value,"1D") ) {
566 *ch_sel = neko_chunks_env();
567 CASE_1D_SEL(LX, *ch_sel);
568 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
569 NEKO_CHUNKS_SEL(LX, *ch_sel));
570 log_message(neko_log_buf);
571 log_end_section();
572 return 1;
573 } else if( !strcmp(env_value,"KSTEP") ) {
574 const int c = neko_eb_env();
575 *eb_sel = c;
576 CASE_KSTEP_SEL(LX, c);
577 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
578 NEKO_EB_SEL(LX, c));
579 log_message(neko_log_buf);
580 log_end_section();
581 return 2;
582 } else if( !strcmp(env_value,"DMMA") ) {
583 if (dmma) {
584 const int c = neko_dmma_env();
585 *nw_sel = c;
586 CASE_DMMA_SEL(LX, c);
587 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
588 NEKO_DMMA_NW(c));
589 log_message(neko_log_buf);
590 log_end_section();
591 return 3;
592 } else {
593 sprintf(neko_log_buf, "DMMA strategy not available for this config");
594 log_error(neko_log_buf);
595 }
596 } else if( !strcmp(env_value,"DMMA_TMA") ) {
597 if (tma) {
598 const int c = neko_dmma_tma_env();
599 *tw_sel = c;
600 CASE_DMMA_TMA_SEL(LX, c);
601 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
602 NEKO_DMMA_NW(c));
603 log_message(neko_log_buf);
604 log_end_section();
605 return 4;
606 } else {
607 sprintf(neko_log_buf,
608 "DMMA_TMA strategy not available for this config");
609 log_error(neko_log_buf);
610 }
611 } else {
612 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
613 log_error(neko_log_buf);
614 }
615 }
616
617 cudaEventCreate(&start);
618 cudaEventCreate(&stop);
619
620 /* Warm every variant before timing anything: each specialisation has to
621 be resident and the clocks at steady state, or whichever variant is
622 timed first is measured on a colder part */
623 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
624 CASE_1D(LX, 0);
625 CASE_1D(LX, 1);
626 CASE_1D(LX, 2);
627 CASE_1D(LX, 3);
628 CASE_KSTEP(LX, 0);
629 if (sweep) {
630 CASE_KSTEP(LX, 1);
631 CASE_KSTEP(LX, 2);
632 }
633 if (dmma) {
634 CASE_DMMA(LX, 0);
635 CASE_DMMA(LX, 1);
636 CASE_DMMA(LX, 2);
637 }
638 if (tma) {
639 CASE_DMMA_TMA(LX, 0);
640 CASE_DMMA_TMA(LX, 1);
641 CASE_DMMA_TMA(LX, 2);
642 }
643 }
644
645 /* Interleaved rounds, best time per variant: timing them one after another
646 in a fixed order lets clock drift bias the comparison by position */
647 for (int r = 0; r < rounds; r++) {
648 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
649 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
650 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
651 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
652 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 0, iters);
653 if (sweep) {
654 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 1, iters);
655 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 2, iters);
656 }
657 if (dmma) {
658 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 0, iters);
659 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 1, iters);
660 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 2, iters);
661 }
662 if (tma) {
663 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 0, iters);
664 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 1, iters);
665 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 2, iters);
666 }
667 }
668
669 NEKO_TUNE_LOG(LX, time1, time2);
670 NEKO_TUNE_LOG_DMMA(LX, time3);
671 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
672
673 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
674 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
675 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
676 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
677 *eb_sel = best;
678 *ch_sel = best1;
679 *nw_sel = best3;
680 *tw_sel = best4;
681
682 if (time1[best1] < time2[best]) {
683 retval = 1;
684 } else {
685 retval = 2;
686 }
687
688 /* The dmma variants join the comparison only where they exist, their
689 candidates are left at NEKO_TUNE_INIT otherwise */
690 float tbest = (retval == 1) ? time1[best1] : time2[best];
691
692 if (time3[best3] < tbest) {
693 retval = 3;
694 tbest = time3[best3];
695 }
696 if (time4[best4] < tbest) {
697 retval = 4;
698 }
699
700 cudaEventDestroy(start);
701 cudaEventDestroy(stop);
702
703 /* The tuner stands in for a real Ax evaluation, and the variants do not
704 sum in the same order, so leave the output of the chosen kernel in w */
705 if (retval == 1) {
706 CASE_1D_SEL(LX, best1);
707 } else if (retval == 2) {
708 CASE_KSTEP_SEL(LX, best);
709 } else if (retval == 3) {
710 CASE_DMMA_SEL(LX, best3);
711 } else {
712 CASE_DMMA_TMA_SEL(LX, best4);
713 }
714
715 if (retval == 1) {
716 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
717 NEKO_CHUNKS_SEL(LX, best1));
718 } else if (retval == 2) {
719 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
720 NEKO_EB_SEL(LX, best));
721 } else if (retval == 3) {
722 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
723 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
724 } else {
725 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
726 NEKO_DMMA_NW(best4));
727 }
728 log_message(neko_log_buf);
729 log_end_section();
730 return retval;
731}
732
733template < const int LX >
734int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
735 void *dxt, void *dyt, void *dzt, void *h1,
736 void *g11, void *g22, void *g33, void *g12,
737 void *g13, void *g23, int *nelv, int *lx,
738 int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel) {
739 cudaEvent_t start, stop;
740 float time1[NEKO_CHUNKS_CANDIDATES];
741 int best1 = 0;
742 float time2[NEKO_EB_CANDIDATES];
743 float time3[NEKO_DMMA_CANDIDATES];
744 int best3 = 0;
745 float time4[NEKO_DMMA_CANDIDATES];
746 int best4 = 0;
747 const int rounds = neko_tune_rounds();
748 const int iters = neko_tune_iters();
749 const int sweep = neko_eb_sweep();
750 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
751 /* Whether the pointers are bulk copy aligned is a property of this call
752 rather than of the kernel, so it is checked rather than assumed, see
753 dmma_tma_aligned() in dmma_tma_kernel.h */
754 const bool tma = dmma_tma_lx_supported<LX>() && cuda_have_tma() &&
755 dmma_tma_aligned(w, u, h1, g11, g22, g33, g12, g13, g23);
756 int best = 0;
757 int retval;
758
759 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
760 time2[c] = NEKO_TUNE_INIT;
761 }
762 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
763 time1[c] = NEKO_TUNE_INIT;
764 }
765 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
766 time3[c] = NEKO_TUNE_INIT;
767 time4[c] = NEKO_TUNE_INIT;
768 }
769
770 const dim3 nblcks_1d((*nelv), 1, 1);
771 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
772
773 char *env_value = NULL;
774 char neko_log_buf[80];
775
776 env_value=getenv("NEKO_AUTOTUNE");
777
778 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
779 log_section(neko_log_buf);
780
781 *eb_sel = 0;
782 *ch_sel = 0;
783 *tw_sel = 0;
784
785 if(env_value) {
786 if( !strcmp(env_value,"1D") ) {
787 *ch_sel = neko_chunks_env();
788 CASE_1D_SEL(LX, *ch_sel);
789 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
790 NEKO_CHUNKS_SEL(LX, *ch_sel));
791 log_message(neko_log_buf);
792 log_end_section();
793 return 1;
794 } else if( !strcmp(env_value,"KSTEP") ) {
795 const int c = neko_eb_env();
796 *eb_sel = c;
797 CASE_KSTEP_PADDED_SEL(LX, c);
798 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
799 NEKO_EB_SEL(LX, c));
800 log_message(neko_log_buf);
801 log_end_section();
802 return 2;
803 } else if( !strcmp(env_value,"DMMA") ) {
804 if (dmma) {
805 const int c = neko_dmma_env();
806 *nw_sel = c;
807 CASE_DMMA_SEL(LX, c);
808 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
809 NEKO_DMMA_NW(c));
810 log_message(neko_log_buf);
811 log_end_section();
812 return 3;
813 } else {
814 sprintf(neko_log_buf, "DMMA strategy not available for this config");
815 log_error(neko_log_buf);
816 }
817 } else if( !strcmp(env_value,"DMMA_TMA") ) {
818 if (tma) {
819 const int c = neko_dmma_tma_env();
820 *tw_sel = c;
821 CASE_DMMA_TMA_SEL(LX, c);
822 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
823 NEKO_DMMA_NW(c));
824 log_message(neko_log_buf);
825 log_end_section();
826 return 4;
827 } else {
828 sprintf(neko_log_buf,
829 "DMMA_TMA strategy not available for this config");
830 log_error(neko_log_buf);
831 }
832 } else {
833 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
834 log_error(neko_log_buf);
835 }
836 }
837
838 cudaEventCreate(&start);
839 cudaEventCreate(&stop);
840
841 /* Warm every variant before timing anything: each specialisation has to
842 be resident and the clocks at steady state, or whichever variant is
843 timed first is measured on a colder part */
844 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
845 CASE_1D(LX, 0);
846 CASE_1D(LX, 1);
847 CASE_1D(LX, 2);
848 CASE_1D(LX, 3);
849 CASE_KSTEP_PADDED(LX, 0);
850 if (sweep) {
851 CASE_KSTEP_PADDED(LX, 1);
852 CASE_KSTEP_PADDED(LX, 2);
853 }
854 if (dmma) {
855 CASE_DMMA(LX, 0);
856 CASE_DMMA(LX, 1);
857 CASE_DMMA(LX, 2);
858 }
859 if (tma) {
860 CASE_DMMA_TMA(LX, 0);
861 CASE_DMMA_TMA(LX, 1);
862 CASE_DMMA_TMA(LX, 2);
863 }
864 }
865
866 /* Interleaved rounds, best time per variant: timing them one after another
867 in a fixed order lets clock drift bias the comparison by position */
868 for (int r = 0; r < rounds; r++) {
869 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
870 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
871 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
872 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
873 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 0, iters);
874 if (sweep) {
875 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 1, iters);
876 NEKO_TUNE_TIME(time2, CASE_KSTEP_PADDED, LX, 2, iters);
877 }
878 if (dmma) {
879 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 0, iters);
880 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 1, iters);
881 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 2, iters);
882 }
883 if (tma) {
884 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 0, iters);
885 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 1, iters);
886 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 2, iters);
887 }
888 }
889
890 NEKO_TUNE_LOG(LX, time1, time2);
891 NEKO_TUNE_LOG_DMMA(LX, time3);
892 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
893
894 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
895 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
896 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
897 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
898 *eb_sel = best;
899 *ch_sel = best1;
900 *nw_sel = best3;
901 *tw_sel = best4;
902
903 if (time1[best1] < time2[best]) {
904 retval = 1;
905 } else {
906 retval = 2;
907 }
908
909 /* The dmma variants join the comparison only where they exist, their
910 candidates are left at NEKO_TUNE_INIT otherwise */
911 float tbest = (retval == 1) ? time1[best1] : time2[best];
912
913 if (time3[best3] < tbest) {
914 retval = 3;
915 tbest = time3[best3];
916 }
917 if (time4[best4] < tbest) {
918 retval = 4;
919 }
920
921 cudaEventDestroy(start);
922 cudaEventDestroy(stop);
923
924 /* Leave the output of the chosen kernel in w, see tune() */
925 if (retval == 1) {
926 CASE_1D_SEL(LX, best1);
927 } else if (retval == 2) {
928 CASE_KSTEP_PADDED_SEL(LX, best);
929 } else if (retval == 3) {
930 CASE_DMMA_SEL(LX, best3);
931 } else {
932 CASE_DMMA_TMA_SEL(LX, best4);
933 }
934
935 if (retval == 1) {
936 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
937 NEKO_CHUNKS_SEL(LX, best1));
938 } else if (retval == 2) {
939 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
940 NEKO_EB_SEL(LX, best));
941 } else if (retval == 3) {
942 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
943 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
944 } else {
945 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
946 NEKO_DMMA_NW(best4));
947 }
948 log_message(neko_log_buf);
949 log_end_section();
950 return retval;
951}
952
953
954/*
955 * Autotuner for the vector Ax
956 *
957 * The vector operator has no 1d variant, so the sweep is the elements per
958 * block candidates of the kstep variant against the warps per block
959 * candidates of the dmma variant -- the same two dimensions the scalar tuner
960 * sweeps, and the same shape as the HIP vector tuner. See the note at the top
961 * of ax_helm_kernel.h for why blocking is not expected to pay here, and why
962 * it is measured rather than assumed. Sampling follows the scalar tuner:
963 * interleaved rounds, best round per candidate.
964 *
965 * @note The section title has to fit the 30 character header field in
966 * log_section(), which truncates rather than wraps -- hence "Ax vector"
967 * rather than the "Ax helm vector" it reads as.
968 */
969template < const int LX >
970int tune_vector(void *au, void *av, void *aw, void *u, void *v, void *w,
971 void *dx, void *dy, void *dz, void *h1,
972 void *g11, void *g22, void *g33, void *g12,
973 void *g13, void *g23, int *nelv, int *lx,
974 int *eb_sel, int *nw_sel, int *tw_sel,
975 int *bw_sel) {
976 cudaEvent_t start, stop;
977 float time2[NEKO_EB_CANDIDATES];
978 int best2 = 0;
979 float time3[NEKO_DMMA_CANDIDATES];
980 int best3 = 0;
981 float time4[NEKO_DMMA_CANDIDATES];
982 int best4 = 0;
983 float time5[NEKO_DMMA_CANDIDATES];
984 int best5 = 0;
985 const int rounds = neko_tune_rounds();
986 const int iters = neko_tune_iters();
987 const int sweep = neko_eb_sweep();
988 const bool dmma = dmma_vector_lx_supported<LX>() && cuda_have_dmma();
989 /* Whether the pointers are bulk copy aligned is a property of this call
990 rather than of the kernel, see dmma_tma_vector_aligned() */
991 const bool tma = dmma_tma_vector_lx_supported<LX>() && cuda_have_tma() &&
992 dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
993 g11, g22, g33, g12, g13, g23);
994 /* Same pointers, and additionally the device has to grant the block the
995 dynamic allocation, see cuda_have_tma_batch() */
996 const bool batch = dmma_tma_batch_lx_supported<LX>() && cuda_have_tma_batch()
997 && dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
998 g11, g22, g33,
999 g12, g13, g23);
1000 int retval;
1001
1002 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1003 time2[c] = NEKO_TUNE_INIT;
1004 }
1005 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
1006 time3[c] = NEKO_TUNE_INIT;
1007 time4[c] = NEKO_TUNE_INIT;
1008 time5[c] = NEKO_TUNE_INIT;
1009 }
1010
1011 const dim3 nblcks_dmma((*nelv), 1, 1);
1012 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
1013
1014 char *env_value = NULL;
1015 char neko_log_buf[80];
1016
1017 env_value=getenv("NEKO_AUTOTUNE");
1018
1019 sprintf(neko_log_buf, "Autotune Ax vector (lx: %d)", *lx);
1020 log_section(neko_log_buf);
1021
1022 *nw_sel = 0;
1023 *tw_sel = 0;
1024 *bw_sel = 0;
1025
1026 if(env_value) {
1027 if( !strcmp(env_value,"KSTEP") || !strcmp(env_value,"1D") ) {
1028 /* No 1d variant here, so either pin lands on kstep */
1029 const int c = neko_eb_env();
1030 *eb_sel = c;
1031 CASE_VECTOR_KSTEP_SEL(LX, c);
1032 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
1033 NEKO_EB_SEL(LX, c));
1034 log_message(neko_log_buf);
1035 log_end_section();
1036 return 2;
1037 } else if( !strcmp(env_value,"DMMA") ) {
1038 if (dmma) {
1039 const int c = neko_dmma_env();
1040 *nw_sel = c;
1041 CASE_VECTOR_DMMA_SEL(LX, c);
1042 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
1043 NEKO_DMMA_NW(c));
1044 log_message(neko_log_buf);
1045 log_end_section();
1046 return 3;
1047 } else {
1048 sprintf(neko_log_buf, "DMMA strategy not available for this config");
1049 log_error(neko_log_buf);
1050 }
1051 } else if( !strcmp(env_value,"DMMA_TMA") ) {
1052 if (tma) {
1053 const int c = neko_dmma_tma_env();
1054 *tw_sel = c;
1055 CASE_VECTOR_DMMA_TMA_SEL(LX, c);
1056 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
1057 NEKO_DMMA_NW(c));
1058 log_message(neko_log_buf);
1059 log_end_section();
1060 return 4;
1061 } else {
1062 sprintf(neko_log_buf,
1063 "DMMA_TMA strategy not available for this config");
1064 log_error(neko_log_buf);
1065 }
1066 } else if( !strcmp(env_value,"DMMA_TMA_BATCH") ) {
1067 if (batch) {
1068 const int c = neko_dmma_tma_env();
1069 *bw_sel = c;
1070 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, c);
1071 sprintf(neko_log_buf,"Set by env : 5 (DMMA_TMA_BATCH, %d warps)",
1072 NEKO_DMMA_NW(c));
1073 log_message(neko_log_buf);
1074 log_end_section();
1075 return 5;
1076 } else {
1077 sprintf(neko_log_buf,
1078 "DMMA_TMA_BATCH strategy not available for this config");
1079 log_error(neko_log_buf);
1080 }
1081 } else {
1082 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1083 log_error(neko_log_buf);
1084 }
1085 }
1086
1087 cudaEventCreate(&start);
1088 cudaEventCreate(&stop);
1089
1090 /* Warm every variant before timing anything, see tune() */
1091 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1092 CASE_VECTOR_KSTEP(LX, 0);
1093 if (sweep) {
1094 CASE_VECTOR_KSTEP(LX, 1);
1095 CASE_VECTOR_KSTEP(LX, 2);
1096 }
1097 if (dmma) {
1098 CASE_VECTOR_DMMA(LX, 0);
1099 CASE_VECTOR_DMMA(LX, 1);
1100 CASE_VECTOR_DMMA(LX, 2);
1101 }
1102 if (tma) {
1103 CASE_VECTOR_DMMA_TMA(LX, 0);
1104 CASE_VECTOR_DMMA_TMA(LX, 1);
1105 CASE_VECTOR_DMMA_TMA(LX, 2);
1106 }
1107 if (batch) {
1108 CASE_VECTOR_DMMA_TMA_BATCH(LX, 0);
1109 CASE_VECTOR_DMMA_TMA_BATCH(LX, 1);
1110 CASE_VECTOR_DMMA_TMA_BATCH(LX, 2);
1111 }
1112 }
1113
1114 for (int r = 0; r < rounds; r++) {
1115 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP, LX, 0, iters);
1116 if (sweep) {
1117 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP, LX, 1, iters);
1118 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP, LX, 2, iters);
1119 }
1120 if (dmma) {
1121 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA, LX, 0, iters);
1122 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA, LX, 1, iters);
1123 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA, LX, 2, iters);
1124 }
1125 if (tma) {
1126 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA, LX, 0, iters);
1127 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA, LX, 1, iters);
1128 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA, LX, 2, iters);
1129 }
1130 if (batch) {
1131 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH, LX, 0, iters);
1132 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH, LX, 1, iters);
1133 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH, LX, 2, iters);
1134 }
1135 }
1136
1137 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1138 if (time2[c] >= NEKO_TUNE_INIT) { continue; }
1139 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call",
1140 NEKO_EB_SEL(LX, c), NEKO_TUNE_US(time2[c], iters));
1141 log_message(neko_log_buf);
1142 }
1143 NEKO_TUNE_LOG_DMMA(LX, time3);
1144 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
1145 NEKO_TUNE_LOG_DMMA_TMA_BATCH(LX, time5);
1146
1147 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1148 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
1149 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
1150 NEKO_TUNE_BEST(time5, best5, NEKO_DMMA_CANDIDATES);
1151 *eb_sel = best2;
1152 *nw_sel = best3;
1153 *tw_sel = best4;
1154 *bw_sel = best5;
1155
1156 /* The dmma variants join the comparison only where they exist, their
1157 candidates are left at NEKO_TUNE_INIT otherwise */
1158 float tbest = time2[best2];
1159
1160 retval = 2;
1161 if (time3[best3] < tbest) { retval = 3; tbest = time3[best3]; }
1162 if (time4[best4] < tbest) { retval = 4; tbest = time4[best4]; }
1163 if (time5[best5] < tbest) { retval = 5; }
1164
1165 cudaEventDestroy(start);
1166 cudaEventDestroy(stop);
1167
1168 /* The tuner stands in for a real Ax evaluation, and the variants do not
1169 sum in the same order, so leave the output of the chosen kernel in
1170 au, av and aw */
1171 if (retval == 2) {
1172 CASE_VECTOR_KSTEP_SEL(LX, best2);
1173 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1174 NEKO_EB_SEL(LX, best2));
1175 } else if (retval == 3) {
1176 CASE_VECTOR_DMMA_SEL(LX, best3);
1177 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
1178 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
1179 } else if (retval == 4) {
1180 CASE_VECTOR_DMMA_TMA_SEL(LX, best4);
1181 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
1182 NEKO_DMMA_NW(best4));
1183 } else {
1184 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, best5);
1185 sprintf(neko_log_buf, "Chose : 5 (DMMA_TMA_BATCH, %d warps)",
1186 NEKO_DMMA_NW(best5));
1187 }
1188 log_message(neko_log_buf);
1189 log_end_section();
1190 return retval;
1191}
1192
1193/* Padded variant of tune_vector(), see the note on bank conflicts above */
1194template < const int LX >
1195int tune_vector_padded(void *au, void *av, void *aw, void *u, void *v, void *w,
1196 void *dx, void *dy, void *dz, void *h1,
1197 void *g11, void *g22, void *g33, void *g12,
1198 void *g13, void *g23, int *nelv, int *lx,
1199 int *eb_sel, int *nw_sel, int *tw_sel,
1200 int *bw_sel) {
1201 cudaEvent_t start, stop;
1202 float time2[NEKO_EB_CANDIDATES];
1203 int best2 = 0;
1204 float time3[NEKO_DMMA_CANDIDATES];
1205 int best3 = 0;
1206 float time4[NEKO_DMMA_CANDIDATES];
1207 int best4 = 0;
1208 float time5[NEKO_DMMA_CANDIDATES];
1209 int best5 = 0;
1210 const int rounds = neko_tune_rounds();
1211 const int iters = neko_tune_iters();
1212 const int sweep = neko_eb_sweep();
1213 const bool dmma = dmma_vector_lx_supported<LX>() && cuda_have_dmma();
1214 /* Whether the pointers are bulk copy aligned is a property of this call
1215 rather than of the kernel, see dmma_tma_vector_aligned() */
1216 const bool tma = dmma_tma_vector_lx_supported<LX>() && cuda_have_tma() &&
1217 dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
1218 g11, g22, g33, g12, g13, g23);
1219 /* Same pointers, and additionally the device has to grant the block the
1220 dynamic allocation, see cuda_have_tma_batch() */
1221 const bool batch = dmma_tma_batch_lx_supported<LX>() && cuda_have_tma_batch()
1222 && dmma_tma_vector_aligned(au, av, aw, u, v, w, h1,
1223 g11, g22, g33,
1224 g12, g13, g23);
1225 int retval;
1226
1227 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1228 time2[c] = NEKO_TUNE_INIT;
1229 }
1230 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
1231 time3[c] = NEKO_TUNE_INIT;
1232 time4[c] = NEKO_TUNE_INIT;
1233 time5[c] = NEKO_TUNE_INIT;
1234 }
1235
1236 const dim3 nblcks_dmma((*nelv), 1, 1);
1237 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
1238
1239 char *env_value = NULL;
1240 char neko_log_buf[80];
1241
1242 env_value=getenv("NEKO_AUTOTUNE");
1243
1244 sprintf(neko_log_buf, "Autotune Ax vector (lx: %d)", *lx);
1245 log_section(neko_log_buf);
1246
1247 *nw_sel = 0;
1248 *tw_sel = 0;
1249 *bw_sel = 0;
1250
1251 if(env_value) {
1252 if( !strcmp(env_value,"KSTEP") || !strcmp(env_value,"1D") ) {
1253 /* No 1d variant here, so either pin lands on kstep */
1254 const int c = neko_eb_env();
1255 *eb_sel = c;
1256 CASE_VECTOR_KSTEP_PADDED_SEL(LX, c);
1257 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
1258 NEKO_EB_SEL(LX, c));
1259 log_message(neko_log_buf);
1260 log_end_section();
1261 return 2;
1262 } else if( !strcmp(env_value,"DMMA") ) {
1263 if (dmma) {
1264 const int c = neko_dmma_env();
1265 *nw_sel = c;
1266 CASE_VECTOR_DMMA_SEL(LX, c);
1267 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
1268 NEKO_DMMA_NW(c));
1269 log_message(neko_log_buf);
1270 log_end_section();
1271 return 3;
1272 } else {
1273 sprintf(neko_log_buf, "DMMA strategy not available for this config");
1274 log_error(neko_log_buf);
1275 }
1276 } else if( !strcmp(env_value,"DMMA_TMA") ) {
1277 if (tma) {
1278 const int c = neko_dmma_tma_env();
1279 *tw_sel = c;
1280 CASE_VECTOR_DMMA_TMA_SEL(LX, c);
1281 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
1282 NEKO_DMMA_NW(c));
1283 log_message(neko_log_buf);
1284 log_end_section();
1285 return 4;
1286 } else {
1287 sprintf(neko_log_buf,
1288 "DMMA_TMA strategy not available for this config");
1289 log_error(neko_log_buf);
1290 }
1291 } else if( !strcmp(env_value,"DMMA_TMA_BATCH") ) {
1292 if (batch) {
1293 const int c = neko_dmma_tma_env();
1294 *bw_sel = c;
1295 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, c);
1296 sprintf(neko_log_buf,"Set by env : 5 (DMMA_TMA_BATCH, %d warps)",
1297 NEKO_DMMA_NW(c));
1298 log_message(neko_log_buf);
1299 log_end_section();
1300 return 5;
1301 } else {
1302 sprintf(neko_log_buf,
1303 "DMMA_TMA_BATCH strategy not available for this config");
1304 log_error(neko_log_buf);
1305 }
1306 } else {
1307 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
1308 log_error(neko_log_buf);
1309 }
1310 }
1311
1312 cudaEventCreate(&start);
1313 cudaEventCreate(&stop);
1314
1315 /* Warm every variant before timing anything, see tune() */
1316 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
1317 CASE_VECTOR_KSTEP_PADDED(LX, 0);
1318 if (sweep) {
1319 CASE_VECTOR_KSTEP_PADDED(LX, 1);
1320 CASE_VECTOR_KSTEP_PADDED(LX, 2);
1321 }
1322 if (dmma) {
1323 CASE_VECTOR_DMMA(LX, 0);
1324 CASE_VECTOR_DMMA(LX, 1);
1325 CASE_VECTOR_DMMA(LX, 2);
1326 }
1327 if (tma) {
1328 CASE_VECTOR_DMMA_TMA(LX, 0);
1329 CASE_VECTOR_DMMA_TMA(LX, 1);
1330 CASE_VECTOR_DMMA_TMA(LX, 2);
1331 }
1332 if (batch) {
1333 CASE_VECTOR_DMMA_TMA_BATCH(LX, 0);
1334 CASE_VECTOR_DMMA_TMA_BATCH(LX, 1);
1335 CASE_VECTOR_DMMA_TMA_BATCH(LX, 2);
1336 }
1337 }
1338
1339 for (int r = 0; r < rounds; r++) {
1340 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_PADDED, LX, 0, iters);
1341 if (sweep) {
1342 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_PADDED, LX, 1, iters);
1343 NEKO_TUNE_TIME(time2, CASE_VECTOR_KSTEP_PADDED, LX, 2, iters);
1344 }
1345 if (dmma) {
1346 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA, LX, 0, iters);
1347 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA, LX, 1, iters);
1348 NEKO_TUNE_TIME(time3, CASE_VECTOR_DMMA, LX, 2, iters);
1349 }
1350 if (tma) {
1351 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA, LX, 0, iters);
1352 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA, LX, 1, iters);
1353 NEKO_TUNE_TIME(time4, CASE_VECTOR_DMMA_TMA, LX, 2, iters);
1354 }
1355 if (batch) {
1356 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH, LX, 0, iters);
1357 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH, LX, 1, iters);
1358 NEKO_TUNE_TIME(time5, CASE_VECTOR_DMMA_TMA_BATCH, LX, 2, iters);
1359 }
1360 }
1361
1362 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
1363 if (time2[c] >= NEKO_TUNE_INIT) { continue; }
1364 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call",
1365 NEKO_EB_SEL(LX, c), NEKO_TUNE_US(time2[c], iters));
1366 log_message(neko_log_buf);
1367 }
1368 NEKO_TUNE_LOG_DMMA(LX, time3);
1369 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
1370 NEKO_TUNE_LOG_DMMA_TMA_BATCH(LX, time5);
1371
1372 NEKO_TUNE_BEST(time2, best2, NEKO_EB_CANDIDATES);
1373 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
1374 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
1375 NEKO_TUNE_BEST(time5, best5, NEKO_DMMA_CANDIDATES);
1376 *eb_sel = best2;
1377 *nw_sel = best3;
1378 *tw_sel = best4;
1379 *bw_sel = best5;
1380
1381 /* The dmma variants join the comparison only where they exist, their
1382 candidates are left at NEKO_TUNE_INIT otherwise */
1383 float tbest = time2[best2];
1384
1385 retval = 2;
1386 if (time3[best3] < tbest) { retval = 3; tbest = time3[best3]; }
1387 if (time4[best4] < tbest) { retval = 4; tbest = time4[best4]; }
1388 if (time5[best5] < tbest) { retval = 5; }
1389
1390 cudaEventDestroy(start);
1391 cudaEventDestroy(stop);
1392
1393 /* Leave the output of the chosen kernel in au, av and aw, see tune() */
1394 if (retval == 2) {
1395 CASE_VECTOR_KSTEP_PADDED_SEL(LX, best2);
1396 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
1397 NEKO_EB_SEL(LX, best2));
1398 } else if (retval == 3) {
1399 CASE_VECTOR_DMMA_SEL(LX, best3);
1400 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
1401 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
1402 } else if (retval == 4) {
1403 CASE_VECTOR_DMMA_TMA_SEL(LX, best4);
1404 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
1405 NEKO_DMMA_NW(best4));
1406 } else {
1407 CASE_VECTOR_DMMA_TMA_BATCH_SEL(LX, best5);
1408 sprintf(neko_log_buf, "Chose : 5 (DMMA_TMA_BATCH, %d warps)",
1409 NEKO_DMMA_NW(best5));
1410 }
1411 log_message(neko_log_buf);
1412 log_end_section();
1413 return retval;
1414}
CASE_DMMA
#define CASE_DMMA(LX, C)
CASE_VECTOR_DMMA
#define CASE_VECTOR_DMMA(LX, C)
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:970
CASE_VECTOR_KSTEP
#define CASE_VECTOR_KSTEP(LX, C)
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:734
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
#define CASE_KSTEP(LX, C)
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_VECTOR_DMMA_TMA_BATCH
#define CASE_VECTOR_DMMA_TMA_BATCH(LX, C)
CASE_LARGE_PADDED
#define CASE_LARGE_PADDED(LX)
CASE_VECTOR_KSTEP_PADDED
#define CASE_VECTOR_KSTEP_PADDED(LX, C)
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_DMMA_TMA
#define CASE_DMMA_TMA(LX, C)
CASE_1D
#define CASE_1D(LX, C)
CASE_VECTOR_KSTEP_SEL
#define CASE_VECTOR_KSTEP_SEL(LX, SEL)
CASE_VECTOR_PADDED
#define CASE_VECTOR_PADDED(LX)
CASE_VECTOR_KSTEP_PADDED_SEL
#define CASE_VECTOR_KSTEP_PADDED_SEL(LX, SEL)
CASE_VECTOR_DMMA_TMA
#define CASE_VECTOR_DMMA_TMA(LX, C)
CASE_KSTEP_PADDED
#define CASE_KSTEP_PADDED(LX, C)
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:1195
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
aw
__global__ void T *__restrict__ T *__restrict__ aw
Definition ax_helm_full_kernel.h:45
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
i
const int i
Definition ax_helm_full_kernel.h:95
dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition ax_helm_full_kernel.h:51
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_full_kernel.h:49
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
av
__global__ void T *__restrict__ av
Definition ax_helm_full_kernel.h:44
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_full_kernel.h:50
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
h1
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1
Definition ax_helm_full_kernel.h:52
g23
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g23
Definition ax_helm_kernel.h:184
g22
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g22
Definition ax_helm_kernel.h:180
g13
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g13
Definition ax_helm_kernel.h:183
nelv
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const int nelv
Definition ax_helm_kernel.h:185
g12
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g12
Definition ax_helm_kernel.h:182
g33
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g33
Definition ax_helm_kernel.h:181
g11
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g11
Definition ax_helm_kernel.h:179
dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
Definition cdtp_kernel.h:115
dzt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
Definition cdtp_kernel.h:116
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition cdtp_kernel.h:114
NEKO_CHUNKS_CANDIDATES
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
NEKO_EB_CANDIDATES
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
NEKO_EB_SEL
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
NEKO_CHUNKS_SEL
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
NEKO_TUNE_TIME
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
Definition elem_block_tune.h:119
neko_eb_env
static int neko_eb_env()
Definition elem_block_tune.h:79
NEKO_TUNE_US
#define NEKO_TUNE_US(T, ITERS)
Definition elem_block_tune.h:139
neko_tune_rounds
static int neko_tune_rounds()
Definition elem_block_tune.h:102
NEKO_TUNE_LOG
#define NEKO_TUNE_LOG(LX, T1, T2)
Definition elem_block_tune.h:149
NEKO_TUNE_INIT
#define NEKO_TUNE_INIT
Definition elem_block_tune.h:65
NEKO_TUNE_BEST
#define NEKO_TUNE_BEST(T, BEST, N)
Definition elem_block_tune.h:141
neko_tune_iters
static int neko_tune_iters()
Definition elem_block_tune.h:110
neko_chunks_env
static int neko_chunks_env()
Definition elem_block_tune.h:91
neko_eb_sweep
static int neko_eb_sweep()
Definition elem_block_tune.h:68
NEKO_TUNE_WARMUP
#define NEKO_TUNE_WARMUP
Definition elem_block_tune.h:66
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:224
neko_dmma_env
static int neko_dmma_env()
Definition dmma_kernel.h:154
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:169
NEKO_DMMA_NW
#define NEKO_DMMA_NW(C)
Definition dmma_kernel.h:124
NEKO_TUNE_LOG_DMMA_TMA_BATCH
#define NEKO_TUNE_LOG_DMMA_TMA_BATCH(LX, T5)
Definition dmma_tma_kernel.h:567
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:580
neko_dmma_tma_env
static int neko_dmma_tma_env()
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)