Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_dudxyz.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 "dudxyz_kernel.h"
39#include "elem_block_tune.h"
41#include <device/cuda/check.h>
42
43extern "C" {
44 #include <common/neko_log.h>
45}
46
47template <const int >
48int tune_dudxyz(void *du, void *u,
49 void *dr, void *ds, void *dt,
50 void *dx, void *dy, void *dz,
51 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel,
52 int *nw_sel, int *tw_sel);
53
54extern "C" {
55
59 void cuda_dudxyz(void *du, void *u,
60 void *dr, void *ds, void *dt,
61 void *dx, void *dy, void *dz,
62 void *jacinv, int *nel, int *lx) {
63
64 static int autotune[16] = { 0 };
65 /* elements per block candidate chosen by the tuner */
66 static int autotune_eb[16] = { 0 };
67 /* chunk candidate chosen for the 1d variant */
68 static int autotune_ch[16] = { 0 };
69 /* warps per block candidate chosen for the dmma variant */
70 static int autotune_nw[16] = { 0 };
71 /* warps per block candidate chosen for the tma staged dmma variant */
72 static int autotune_tw[16] = { 0 };
73
74 const dim3 nthrds_1d(1024, 1, 1);
75 const dim3 nthrds_kstep((*lx), (*lx), 1);
76 const dim3 nblcks((*nel), 1, 1);
77 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
78
79#define CASE_1D(LX, C) \
80 dudxyz_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
81 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
82 ((real *) du, (real *) u, \
83 (real *) dr, (real *) ds, (real *) dt, \
84 (real *) dx, (real *) dy, (real *) dz, \
85 (real *) jacinv); \
86 CUDA_CHECK(cudaGetLastError());
87
88/* Runtime dispatch onto the tuned chunk candidate */
89#define CASE_1D_SEL(LX, SEL) \
90 switch (SEL) { \
91 case 0: CASE_1D(LX, 0); break; \
92 case 1: CASE_1D(LX, 1); break; \
93 case 2: CASE_1D(LX, 2); break; \
94 default: CASE_1D(LX, 3); break; \
95 }
96
97#define CASE_KSTEP(LX, C) \
98 dudxyz_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
99 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
100 ((real *) du, (real *) u, \
101 (real *) dr, (real *) ds, (real *) dt, \
102 (real *) dx, (real *) dy, (real *) dz, \
103 (real *) jacinv, *nel); \
104 CUDA_CHECK(cudaGetLastError());
105
106/* Runtime dispatch onto the tuned candidate */
107#define CASE_KSTEP_SEL(LX, SEL) \
108 switch (SEL) { \
109 case 0: CASE_KSTEP(LX, 0); break; \
110 case 1: CASE_KSTEP(LX, 1); break; \
111 default: CASE_KSTEP(LX, 2); break; \
112 }
113
114#define CASE_DMMA(LX, C) \
115 dudxyz_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
116 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
117 ((real *) du, (real *) u, \
118 (real *) dr, (real *) ds, (real *) dt, \
119 (real *) dx, (real *) dy, (real *) dz, \
120 (real *) jacinv, *nel); \
121 CUDA_CHECK(cudaGetLastError());
122
123/* Runtime dispatch onto the tuned warps per block candidate */
124#define CASE_DMMA_SEL(LX, SEL) \
125 switch (SEL) { \
126 case 0: CASE_DMMA(LX, 0); break; \
127 case 1: CASE_DMMA(LX, 1); break; \
128 default: CASE_DMMA(LX, 2); break; \
129 }
130
131/*
132 * The TMA staged dmma variant. Same grid and the same warps per block
133 * candidates as CASE_DMMA -- at the one lx it supports NEKO_DMMA_PACK is 1,
134 * so the grid is nel blocks -- but it takes no nel: there is no packed tail
135 * to clamp when a block is exactly an element. See dmma_tma_kernel.h.
136 */
137#define CASE_DMMA_TMA(LX, C) \
138 dudxyz_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
139 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
140 ((real *) du, (real *) u, \
141 (real *) dr, (real *) ds, (real *) dt, \
142 (real *) dx, (real *) dy, (real *) dz, \
143 (real *) jacinv); \
144 CUDA_CHECK(cudaGetLastError());
145
146/* Runtime dispatch onto the tuned warps per block candidate */
147#define CASE_DMMA_TMA_SEL(LX, SEL) \
148 switch (SEL) { \
149 case 0: CASE_DMMA_TMA(LX, 0); break; \
150 case 1: CASE_DMMA_TMA(LX, 1); break; \
151 default: CASE_DMMA_TMA(LX, 2); break; \
152 }
153
154 #define CASE(LX) \
155 case LX: \
156 if(autotune[LX] == 0 ) { \
157 autotune[LX]=tune_dudxyz<LX>(du, u, \
158 dr, ds, dt, \
159 dx, dy, dz, \
160 jacinv, nel, lx, &autotune_eb[LX], \
161 &autotune_ch[LX], &autotune_nw[LX], \
162 &autotune_tw[LX]); \
163 } else if (autotune[LX] == 1 ) { \
164 CASE_1D_SEL(LX, autotune_ch[LX]); \
165 } else if (autotune[LX] == 2 ) { \
166 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
167 } else if (autotune[LX] == 3 ) { \
168 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
169 } else if (autotune[LX] == 4 ) { \
170 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
171 } \
172 break
173
174#define CASE_LARGE(LX) \
175 case LX: \
176 CASE_KSTEP(LX, 0); \
177 break
178
179
180 if ((*lx) < 11) {
181 switch(*lx) {
182 CASE(2);
183 CASE(3);
184 CASE(4);
185 CASE(5);
186 CASE(6);
187 CASE(7);
188 CASE(8);
189 CASE(9);
190 CASE(10);
191 default:
192 {
193 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
194 exit(1);
195 }
196 }
197 }
198 else {
199 switch(*lx) {
200 CASE_LARGE(11);
201 CASE_LARGE(12);
202 CASE_LARGE(13);
203 CASE_LARGE(14);
204 CASE_LARGE(15);
205 CASE_LARGE(16);
206 default:
207 {
208 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
209 exit(1);
210 }
211 }
212 }
213 }
214}
215
216template < const int LX >
217int tune_dudxyz(void *du, void *u,
218 void *dr, void *ds, void *dt,
219 void *dx, void *dy, void *dz,
220 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel,
221 int *nw_sel, int *tw_sel) {
224 int best1 = 0;
227 int best3 = 0;
229 int best4 = 0;
230 const int rounds = neko_tune_rounds();
231 const int iters = neko_tune_iters();
232 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
233 elements per block sweep is off */
234 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
235 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
236 const int ch_pin = neko_chunks_pin();
237 const int eb_pin = neko_eb_pin();
238 const int nw_pin = neko_dmma_pin();
239 const int tw_pin = neko_dmma_tma_pin();
240 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
241 int strat = 0;
242 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
243 /* Whether the pointers are bulk copy aligned is a property of this call
244 rather than of the kernel, so it is checked rather than assumed, see
245 dmma_tma_dudxyz_aligned() in dmma_tma_kernel.h */
248 int best = 0;
249 int retval;
250
251 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
253 }
254 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
256 }
257 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
260 }
261
262 const dim3 nthrds_1d(1024, 1, 1);
263 const dim3 nthrds_kstep((*lx), (*lx), 1);
264 const dim3 nblcks((*nel), 1, 1);
265 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
266
267 char *env_value = NULL;
268 char neko_log_buf[80];
269
270 env_value=getenv("NEKO_AUTOTUNE");
271
272 sprintf(neko_log_buf, "Autotune dudxyz (lx: %d)", *lx);
274
275 *eb_sel = 0;
276 *ch_sel = 0;
277 *nw_sel = 0;
278 *tw_sel = 0;
279
280 /*
281 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
282 * below is narrowed to that one kernel family, but its geometry -- the
283 * chunk size, the elements per block, the warps per block -- is still
284 * measured candidate against candidate. A formulation this build or this
285 * device does not have is reported and ignored, leaving the full sweep.
286 */
287 if(env_value) {
288 if( !strcmp(env_value,"1D") ) {
289 strat = 1;
290 } else if( !strcmp(env_value,"KSTEP") ) {
291 strat = 2;
292 } else if( !strcmp(env_value,"DMMA") ) {
293 if (dmma) {
294 strat = 3;
295 } else {
296 sprintf(neko_log_buf, "DMMA strategy not available for this config");
298 }
299 } else if( !strcmp(env_value,"DMMA_TMA") ) {
300 if (tma) {
301 strat = 4;
302 } else {
304 "DMMA_TMA strategy not available for this config");
306 }
307 } else {
308 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
310 }
311 }
312
313 /* Geometry of the pinned formulation, if its own variable fixes that too.
314 Both pinned leaves nothing to measure, so the kernel is launched once and
315 reported, which is what pinning has always done */
316 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
317 (strat == 3) ? nw_pin : (strat == 4) ? tw_pin : -1;
318
319 if (pin >= 0) {
320 switch (strat) {
321 case 1:
322 *ch_sel = pin;
324 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
326 break;
327 case 2:
328 *eb_sel = pin;
330 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
331 NEKO_EB_SEL(LX, pin));
332 break;
333 case 3:
334 *nw_sel = pin;
336 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
338 break;
339 default:
340 *tw_sel = pin;
342 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
344 break;
345 }
348 return strat;
349 }
350
351 if (strat) {
352 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
354 }
355
356 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
357 const bool try_1d = (strat == 0 || strat == 1);
358 const bool try_kstep = (strat == 0 || strat == 2);
359 const bool try_dmma = dmma && (strat == 0 || strat == 3);
360 const bool try_tma = tma && (strat == 0 || strat == 4);
361
364
365 /* Warm every variant before timing anything: each specialisation has to be
366 resident and the clocks at steady state, or whichever is timed first is
367 measured on a colder part */
368 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
370 CASE_1D_SEL(LX, c);
371 }
373 CASE_KSTEP_SEL(LX, c);
374 }
376 CASE_DMMA_SEL(LX, c);
377 }
380 }
381 }
382
383 /* Interleaved rounds, best time per variant */
384 for (int r = 0; r < rounds; r++) {
387 }
390 }
393 }
396 }
397 }
398
402
407 *eb_sel = best;
408 *ch_sel = best1;
409 *nw_sel = best3;
410 *tw_sel = best4;
411
412 if (time1[best1] < time2[best]) {
413 retval = 1;
414 } else {
415 retval = 2;
416 }
417
418 /* The dmma variants join the comparison only where they exist, their
419 candidates are left at NEKO_TUNE_INIT otherwise */
420 float tbest = (retval == 1) ? time1[best1] : time2[best];
421
422 if (time3[best3] < tbest) {
423 retval = 3;
424 tbest = time3[best3];
425 }
426 if (time4[best4] < tbest) {
427 retval = 4;
428 }
429
430 /* Leave the chosen kernel's output in place: the tuner stands in for a real
431 evaluation and the variants do not sum in the same order */
432 if (retval == 1) {
434 } else if (retval == 2) {
436 } else if (retval == 3) {
438 } else {
440 }
441
442 if (retval == 1) {
443 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
445 } else if (retval == 2) {
446 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
448 } else if (retval == 3) {
449 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
451 } else {
452 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
454 }
457 return retval;
458}
__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)
__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__ 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__ jacinv
const int i
__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
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
__global__ void const T *__restrict__ const T *__restrict__ dr
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
static int neko_chunks_pin()
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
static int neko_tune_rounds()
static int neko_eb_pin()
#define NEKO_TUNE_LOG(LX, T1, T2)
#define NEKO_TUNE_INIT
#define NEKO_TUNE_BEST(T, BEST, N)
static int neko_tune_iters()
#define NEKO_TUNE_FOR(C, ON, PIN, N)
static int neko_eb_sweep()
#define NEKO_TUNE_WARMUP
#define NEKO_DMMA_CANDIDATES
static bool cuda_have_dmma()
#define NEKO_DMMA_PACK(LX)
#define NEKO_TUNE_LOG_DMMA(LX, T3)
#define NEKO_DMMA_NW(C)
static int neko_dmma_pin()
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
static int neko_dmma_tma_pin()
static bool dmma_tma_dudxyz_aligned(const void *du, const void *u, const void *dr, const void *ds, const void *dt, const void *jacinv)
static bool cuda_have_tma()
void log_error(char *msg)
void log_message(char *msg)
void log_end_section()
void log_section(char *msg)
void cuda_dudxyz(void *du, void *u, void *dr, void *ds, void *dt, void *dx, void *dy, void *dz, void *jacinv, int *nel, int *lx)
Definition opr_dudxyz.cu:59
#define CASE_DMMA_SEL(LX, SEL)
#define CASE_KSTEP_SEL(LX, SEL)
#define CASE(LX)
#define CASE_1D_SEL(LX, SEL)
int tune_dudxyz(void *du, void *u, void *dr, void *ds, void *dt, void *dx, void *dy, void *dz, void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
#define CASE_LARGE(LX)
#define CASE_DMMA_TMA_SEL(LX, SEL)