Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_opgrad.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 "opgrad_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_opgrad(void *ux, void *uy, void *uz, void *u,
49 void *dx, void *dy, void *dz,
50 void *drdx, void *dsdx, void *dtdx,
51 void *drdy, void *dsdy, void *dtdy,
52 void *drdz, void *dsdz, void *dtdz,
53 void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel,
54 int *nw_sel, int *tw_sel);
55
56extern "C" {
57
61 void cuda_opgrad(void *ux, void *uy, void *uz, void *u,
62 void *dx, void *dy, void *dz,
63 void *drdx, void *dsdx, void *dtdx,
64 void *drdy, void *dsdy, void *dtdy,
65 void *drdz, void *dsdz, void *dtdz,
66 void *w3, int *nel, int *lx) {
67
68 static int autotune[17] = { 0 };
69 /* elements per block candidate chosen by the tuner */
70 static int autotune_eb[17] = { 0 };
71 /* chunk candidate chosen for the 1d variant */
72 static int autotune_ch[17] = { 0 };
73 /* warps per block candidate chosen for the dmma variant */
74 static int autotune_nw[17] = { 0 };
75 /* warps per block candidate chosen for the tma staged dmma variant */
76 static int autotune_tw[17] = { 0 };
77
78 const dim3 nthrds_1d(1024, 1, 1);
79 const dim3 nthrds_kstep((*lx), (*lx), 1);
80 const dim3 nblcks((*nel), 1, 1);
81 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
82
83#define CASE_1D(LX, C) \
84 opgrad_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
85 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
86 ((real *) ux, (real *) uy, (real *) uz, (real *) u, \
87 (real *) dx, (real *) dy, (real *) dz, \
88 (real *) drdx, (real *) dsdx, (real *) dtdx, \
89 (real *) drdy, (real *) dsdy, (real *) dtdy, \
90 (real *) drdz, (real *) dsdz, (real *) dtdz, \
91 (real *) w3); \
92 CUDA_CHECK(cudaGetLastError());
93
94/* Runtime dispatch onto the tuned chunk candidate */
95#define CASE_1D_SEL(LX, SEL) \
96 switch (SEL) { \
97 case 0: CASE_1D(LX, 0); break; \
98 case 1: CASE_1D(LX, 1); break; \
99 case 2: CASE_1D(LX, 2); break; \
100 default: CASE_1D(LX, 3); break; \
101 }
102
103
104#define CASE_KSTEP(LX, C) \
105 opgrad_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
106 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
107 ((real *) ux, (real *) uy, (real *) uz, (real *) u, \
108 (real *) dx, (real *) dy, (real *) dz, \
109 (real *) drdx, (real *) dsdx, (real *) dtdx, \
110 (real *) drdy, (real *) dsdy, (real *) dtdy, \
111 (real *) drdz, (real *) dsdz, (real *) dtdz, \
112 (real *) w3, *nel); \
113 CUDA_CHECK(cudaGetLastError());
114
115/* Runtime dispatch onto the tuned candidate */
116#define CASE_KSTEP_SEL(LX, SEL) \
117 switch (SEL) { \
118 case 0: CASE_KSTEP(LX, 0); break; \
119 case 1: CASE_KSTEP(LX, 1); break; \
120 default: CASE_KSTEP(LX, 2); break; \
121 }
122
123#define CASE_DMMA(LX, C) \
124 opgrad_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
125 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
126 ((real *) ux, (real *) uy, (real *) uz, (real *) u, \
127 (real *) dx, (real *) dy, (real *) dz, \
128 (real *) drdx, (real *) dsdx, (real *) dtdx, \
129 (real *) drdy, (real *) dsdy, (real *) dtdy, \
130 (real *) drdz, (real *) dsdz, (real *) dtdz, \
131 (real *) w3, *nel); \
132 CUDA_CHECK(cudaGetLastError());
133
134/* Runtime dispatch onto the tuned warps per block candidate */
135#define CASE_DMMA_SEL(LX, SEL) \
136 switch (SEL) { \
137 case 0: CASE_DMMA(LX, 0); break; \
138 case 1: CASE_DMMA(LX, 1); break; \
139 default: CASE_DMMA(LX, 2); break; \
140 }
141
142/*
143 * The TMA staged dmma variant. Same grid and the same warps per block
144 * candidates as CASE_DMMA -- at the one lx it supports NEKO_DMMA_PACK is 1,
145 * so the grid is nel blocks -- but unlike every other launch here it carries
146 * a dynamic shared memory allocation, 54800 B, past the 48 kB a block gets for
147 * free. The kernel has to opt into that once before its first launch; see
148 * opgrad_dmma_tma_optin(), which is a predictable branch after that.
149 */
150#define CASE_DMMA_TMA(LX, C) \
151 (void) opgrad_dmma_tma_optin<real, LX, NEKO_DMMA_NW(C)>(); \
152 opgrad_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
153 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), \
154 NEKO_OPGRAD_TMA_SMEM, stream>>> \
155 ((real *) ux, (real *) uy, (real *) uz, (real *) u, \
156 (real *) dx, (real *) dy, (real *) dz, \
157 (real *) drdx, (real *) dsdx, (real *) dtdx, \
158 (real *) drdy, (real *) dsdy, (real *) dtdy, \
159 (real *) drdz, (real *) dsdz, (real *) dtdz, \
160 (real *) w3); \
161 CUDA_CHECK(cudaGetLastError());
162
163/* Runtime dispatch onto the tuned warps per block candidate */
164#define CASE_DMMA_TMA_SEL(LX, SEL) \
165 switch (SEL) { \
166 case 0: CASE_DMMA_TMA(LX, 0); break; \
167 case 1: CASE_DMMA_TMA(LX, 1); break; \
168 default: CASE_DMMA_TMA(LX, 2); break; \
169 }
170
171#define CASE(LX) \
172 case LX: \
173 if(autotune[LX] == 0 ) { \
174 autotune[LX]=tune_opgrad<LX>(ux, uy, uz, u, \
175 dx, dy, dz, \
176 drdx, dsdx, dtdx, \
177 drdy, dsdy, dtdy, \
178 drdz, dsdz, dtdz, \
179 w3, nel, lx, &autotune_eb[LX], \
180 &autotune_ch[LX], &autotune_nw[LX], \
181 &autotune_tw[LX]); \
182 } else if (autotune[LX] == 1 ) { \
183 CASE_1D_SEL(LX, autotune_ch[LX]); \
184 } else if (autotune[LX] == 2 ) { \
185 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
186 } else if (autotune[LX] == 3 ) { \
187 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
188 } else if (autotune[LX] == 4 ) { \
189 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
190 } \
191 break
192
193 switch(*lx) {
194 CASE(2);
195 CASE(3);
196 CASE(4);
197 CASE(5);
198 CASE(6);
199 CASE(7);
200 CASE(8);
201 CASE(9);
202 CASE(10);
203 CASE(11);
204 CASE(12);
205 CASE(13);
206 CASE(14);
207 CASE(15);
208 CASE(16);
209 default:
210 {
211 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
212 exit(1);
213 }
214 }
215 }
216}
217
218template < const int LX >
219int tune_opgrad(void *ux, void *uy, void *uz, void *u,
220 void *dx, void *dy, void *dz,
221 void *drdx, void *dsdx, void *dtdx,
222 void *drdy, void *dsdy, void *dtdy,
223 void *drdz, void *dsdz, void *dtdz,
224 void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel,
225 int *nw_sel, int *tw_sel) {
228 int best1 = 0;
231 int best3 = 0;
233 int best4 = 0;
234 const int rounds = neko_tune_rounds();
235 const int iters = neko_tune_iters();
236 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
237 elements per block sweep is off */
238 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
239 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
240 const int ch_pin = neko_chunks_pin();
241 const int eb_pin = neko_eb_pin();
242 const int nw_pin = neko_dmma_pin();
243 const int tw_pin = neko_dmma_tma_pin();
244 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
245 int strat = 0;
246 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
247 /* Whether the pointers are bulk copy aligned is a property of this call
248 rather than of the kernel, so it is checked rather than assumed, see
249 dmma_tma_kernel.h. The shared memory gate is a device query too: this
250 variant asks for more than a block gets by default */
251 const bool tma = dmma_tma_opgrad_lx_supported<LX>() &&
254 drdy, dsdy, dtdy,
255 drdz, dsdz, dtdz);
256 int best = 0;
257 int retval;
258
259 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
261 }
262 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
264 }
265 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
268 }
269
270 const dim3 nthrds_1d(1024, 1, 1);
271 const dim3 nthrds_kstep((*lx), (*lx), 1);
272 const dim3 nblcks((*nel), 1, 1);
273 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
274
275 char *env_value = NULL;
276 char neko_log_buf[80];
277
278 env_value=getenv("NEKO_AUTOTUNE");
279
280 sprintf(neko_log_buf, "Autotune opgrad (lx: %d)", *lx);
282
283 *eb_sel = 0;
284 *ch_sel = 0;
285 *nw_sel = 0;
286 *tw_sel = 0;
287
288 /*
289 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
290 * below is narrowed to that one kernel family, but its geometry -- the
291 * chunk size, the elements per block, the warps per block -- is still
292 * measured candidate against candidate. A formulation this build or this
293 * device does not have is reported and ignored, leaving the full sweep.
294 */
295 if(env_value) {
296 if( !strcmp(env_value,"1D") ) {
297 strat = 1;
298 } else if( !strcmp(env_value,"KSTEP") ) {
299 strat = 2;
300 } else if( !strcmp(env_value,"DMMA") ) {
301 if (dmma) {
302 strat = 3;
303 } else {
304 sprintf(neko_log_buf, "DMMA strategy not available for this config");
306 }
307 } else if( !strcmp(env_value,"DMMA_TMA") ) {
308 if (tma) {
309 strat = 4;
310 } else {
312 "DMMA_TMA strategy not available for this config");
314 }
315 } else {
316 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
318 }
319 }
320
321 /* Geometry of the pinned formulation, if its own variable fixes that too.
322 Both pinned leaves nothing to measure, so the kernel is launched once and
323 reported, which is what pinning has always done */
324 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
325 (strat == 3) ? nw_pin : (strat == 4) ? tw_pin : -1;
326
327 if (pin >= 0) {
328 switch (strat) {
329 case 1:
330 *ch_sel = pin;
332 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
334 break;
335 case 2:
336 *eb_sel = pin;
338 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
339 NEKO_EB_SEL(LX, pin));
340 break;
341 case 3:
342 *nw_sel = pin;
344 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
346 break;
347 default:
348 *tw_sel = pin;
350 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
352 break;
353 }
356 return strat;
357 }
358
359 if (strat) {
360 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
362 }
363
364 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
365 const bool try_1d = (strat == 0 || strat == 1);
366 const bool try_kstep = (strat == 0 || strat == 2);
367 const bool try_dmma = dmma && (strat == 0 || strat == 3);
368 const bool try_tma = tma && (strat == 0 || strat == 4);
369
372
373 /* Warm every variant before timing anything: each specialisation has to be
374 resident and the clocks at steady state, or whichever is timed first is
375 measured on a colder part */
376 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
378 CASE_1D_SEL(LX, c);
379 }
381 CASE_KSTEP_SEL(LX, c);
382 }
384 CASE_DMMA_SEL(LX, c);
385 }
388 }
389 }
390
391 /* Interleaved rounds, best time per variant */
392 for (int r = 0; r < rounds; r++) {
395 }
398 }
401 }
404 }
405 }
406
410
415 *eb_sel = best;
416 *ch_sel = best1;
417 *nw_sel = best3;
418 *tw_sel = best4;
419
420 if (time1[best1] < time2[best]) {
421 retval = 1;
422 } else {
423 retval = 2;
424 }
425
426 /* The dmma variants join the comparison only where they exist, their
427 candidates are left at NEKO_TUNE_INIT otherwise */
428 float tbest = (retval == 1) ? time1[best1] : time2[best];
429
430 if (time3[best3] < tbest) {
431 retval = 3;
432 tbest = time3[best3];
433 }
434 if (time4[best4] < tbest) {
435 retval = 4;
436 }
437
438 /* Leave the chosen kernel's output in place: the tuner stands in for a real
439 evaluation and the variants do not sum in the same order */
440 if (retval == 1) {
442 } else if (retval == 2) {
444 } else if (retval == 3) {
446 } else {
448 }
449
450 if (retval == 1) {
451 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
453 } else if (retval == 2) {
454 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
456 } else if (retval == 3) {
457 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
459 } else {
460 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
462 }
465 return retval;
466}
__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__ dtdy
__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__ dtdx
__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__ dtdz
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__ 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__ dsdz
__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 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__ drdz
__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__ drdx
__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__ dsdx
__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__ dsdy
__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__ drdy
__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__ w3
#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
__global__ void T *__restrict__ uy
__global__ void T *__restrict__ T *__restrict__ uz
#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 bool cuda_have_tma_opgrad()
static int neko_dmma_tma_pin()
static bool dmma_tma_opgrad_aligned(const void *u, const void *drdx, const void *dsdx, const void *dtdx, const void *drdy, const void *dsdy, const void *dtdy, const void *drdz, const void *dsdz, const void *dtdz)
void log_error(char *msg)
void log_message(char *msg)
void log_end_section()
void log_section(char *msg)
#define CASE_DMMA_SEL(LX, SEL)
#define CASE_KSTEP_SEL(LX, SEL)
#define CASE(LX)
#define CASE_1D_SEL(LX, SEL)
int tune_opgrad(void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
void cuda_opgrad(void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx)
Definition opr_opgrad.cu:61
#define CASE_DMMA_TMA_SEL(LX, SEL)