Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_cdtp.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 "cdtp_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_cdtp(void *dtx, void *x,
49 void *dr, void *ds, void *dt,
50 void *dxt, void *dyt, void *dzt,
51 void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel,
52 int *nw_sel, int *tw_sel);
53
54extern "C" {
55
59 void cuda_cdtp(void *dtx, void *x,
60 void *dr, void *ds, void *dt,
61 void *dxt, void *dyt, void *dzt,
62 void *w3, int *nel, int *lx) {
63
64 static int autotune[17] = { 0 };
65 /* elements per block candidate chosen by the tuner */
66 static int autotune_eb[17] = { 0 };
67 /* chunk candidate chosen for the 1d variant */
68 static int autotune_ch[17] = { 0 };
69 /* warps per block candidate chosen for the dmma variant */
70 static int autotune_nw[17] = { 0 };
71 /* warps per block candidate chosen for the tma staged dmma variant */
72 static int autotune_tw[17] = { 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 cdtp_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
81 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
82 ((real *) dtx, (real *) x, \
83 (real *) dr, (real *) ds, (real *) dt, \
84 (real *) dxt, (real *) dyt, (real *) dzt, \
85 (real *) w3); \
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 cdtp_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
99 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
100 ((real *) dtx, (real *) x, \
101 (real *) dr, (real *) ds, (real *) dt, \
102 (real *) dxt, (real *) dyt, (real *) dzt, \
103 (real *) w3, *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 cdtp_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
116 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
117 ((real *) dtx, (real *) x, \
118 (real *) dr, (real *) ds, (real *) dt, \
119 (real *) dxt, (real *) dyt, (real *) dzt, \
120 (real *) w3, *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. Its block is 22032 B, inside the
136 * 48 kB a block gets for free, so unlike the opgrad and conv1 variants it
137 * needs no dynamic allocation and no opt-in. See dmma_tma_kernel.h.
138 */
139#define CASE_DMMA_TMA(LX, C) \
140 cdtp_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
141 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
142 ((real *) dtx, (real *) x, \
143 (real *) dr, (real *) ds, (real *) dt, \
144 (real *) dxt, (real *) dyt, (real *) dzt, \
145 (real *) w3); \
146 CUDA_CHECK(cudaGetLastError());
147
148/* Runtime dispatch onto the tuned warps per block candidate */
149#define CASE_DMMA_TMA_SEL(LX, SEL) \
150 switch (SEL) { \
151 case 0: CASE_DMMA_TMA(LX, 0); break; \
152 case 1: CASE_DMMA_TMA(LX, 1); break; \
153 default: CASE_DMMA_TMA(LX, 2); break; \
154 }
155
156#define CASE(LX) \
157 case LX: \
158 if(autotune[LX] == 0 ) { \
159 autotune[LX]=tune_cdtp<LX>(dtx, x, \
160 dr, ds, dt, \
161 dxt, dyt, dzt, \
162 w3, nel, lx, &autotune_eb[LX], \
163 &autotune_ch[LX], &autotune_nw[LX], \
164 &autotune_tw[LX]); \
165 } else if (autotune[LX] == 1 ) { \
166 CASE_1D_SEL(LX, autotune_ch[LX]); \
167 } else if (autotune[LX] == 2 ) { \
168 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
169 } else if (autotune[LX] == 3 ) { \
170 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
171 } else if (autotune[LX] == 4 ) { \
172 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
173 } \
174 break
175
176#define CASE_LARGE(LX) \
177 case LX: \
178 CASE_KSTEP(LX, 0); \
179 break
180
181
182 if ((*lx) < 13) {
183 switch(*lx) {
184 CASE(2);
185 CASE(3);
186 CASE(4);
187 CASE(5);
188 CASE(6);
189 CASE(7);
190 CASE(8);
191 CASE(9);
192 CASE(10);
193 CASE(11);
194 CASE(12);
195 default:
196 {
197 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
198 exit(1);
199 }
200 }
201 }
202 else {
203 switch(*lx) {
204 CASE_LARGE(13);
205 CASE_LARGE(14);
206 CASE_LARGE(15);
207 CASE_LARGE(16);
208 default:
209 {
210 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
211 exit(1);
212 }
213 }
214 }
215 }
216}
217
218template < const int LX >
219int tune_cdtp(void *dtx, void *x,
220 void *dr, void *ds, void *dt,
221 void *dxt, void *dyt, void *dzt,
222 void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel,
223 int *nw_sel, int *tw_sel) {
226 int best1 = 0;
229 int best3 = 0;
231 int best4 = 0;
232 const int rounds = neko_tune_rounds();
233 const int iters = neko_tune_iters();
234 const int sweep = neko_eb_sweep();
235 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
236 /* Whether the pointers are bulk copy aligned is a property of this call
237 rather than of the kernel, so it is checked rather than assumed, see
238 dmma_tma_kernel.h. The shared memory gate is a device query too: this
239 variant asks for more than a block gets by default */
242 int best = 0;
243 int retval;
244
245 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
247 }
248 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
250 }
251 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
254 }
255
256 const dim3 nthrds_1d(1024, 1, 1);
257 const dim3 nthrds_kstep((*lx), (*lx), 1);
258 const dim3 nblcks((*nel), 1, 1);
259 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
260
261 char *env_value = NULL;
262 char neko_log_buf[80];
263
264 env_value=getenv("NEKO_AUTOTUNE");
265
266 sprintf(neko_log_buf, "Autotune cdtp (lx: %d)", *lx);
268
269 *eb_sel = 0;
270 *ch_sel = 0;
271 *nw_sel = 0;
272 *tw_sel = 0;
273
274 if(env_value) {
275 if( !strcmp(env_value,"1D") ) {
278 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
282 return 1;
283 } else if( !strcmp(env_value,"KSTEP") ) {
284 *eb_sel = neko_eb_env();
286 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
290 return 2;
291 } else if( !strcmp(env_value,"DMMA") ) {
292 if (dmma) {
293 const int c = neko_dmma_env();
294 *nw_sel = c;
295 CASE_DMMA_SEL(LX, c);
296 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
297 NEKO_DMMA_NW(c));
300 return 3;
301 } else {
302 sprintf(neko_log_buf, "DMMA strategy not available for this config");
304 }
305 } else if( !strcmp(env_value,"DMMA_TMA") ) {
306 if (tma) {
307 const int c = neko_dmma_tma_env();
308 *tw_sel = c;
310 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
311 NEKO_DMMA_NW(c));
314 return 4;
315 } else {
317 "DMMA_TMA strategy not available for this config");
319 }
320 } else {
321 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
323 }
324 }
325
328
329 /* Warm every variant before timing anything: each specialisation has to be
330 resident and the clocks at steady state, or whichever is timed first is
331 measured on a colder part */
332 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
333 CASE_1D(LX, 0);
334 CASE_1D(LX, 1);
335 CASE_1D(LX, 2);
336 CASE_1D(LX, 3);
337 CASE_KSTEP(LX, 0);
338 if (sweep) {
339 CASE_KSTEP(LX, 1);
340 CASE_KSTEP(LX, 2);
341 }
342 if (dmma) {
343 CASE_DMMA(LX, 0);
344 CASE_DMMA(LX, 1);
345 CASE_DMMA(LX, 2);
346 }
347 if (tma) {
348 CASE_DMMA_TMA(LX, 0);
349 CASE_DMMA_TMA(LX, 1);
350 CASE_DMMA_TMA(LX, 2);
351 }
352 }
353
354 /* Interleaved rounds, best time per variant */
355 for (int r = 0; r < rounds; r++) {
361 if (sweep) {
364 }
365 if (dmma) {
369 }
370 if (tma) {
374 }
375 }
376
380
385 *eb_sel = best;
386 *ch_sel = best1;
387 *nw_sel = best3;
388 *tw_sel = best4;
389
390 if (time1[best1] < time2[best]) {
391 retval = 1;
392 } else {
393 retval = 2;
394 }
395
396 /* The dmma variants join the comparison only where they exist, their
397 candidates are left at NEKO_TUNE_INIT otherwise */
398 float tbest = (retval == 1) ? time1[best1] : time2[best];
399
400 if (time3[best3] < tbest) {
401 retval = 3;
402 tbest = time3[best3];
403 }
404 if (time4[best4] < tbest) {
405 retval = 4;
406 }
407
408 /* Leave the chosen kernel's output in place: the tuner stands in for a real
409 evaluation and the variants do not sum in the same order */
410 if (retval == 1) {
412 } else if (retval == 2) {
414 } else if (retval == 3) {
416 } else {
418 }
419
420 if (retval == 1) {
421 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
423 } else if (retval == 2) {
424 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
426 } else if (retval == 3) {
427 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
429 } else {
430 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
432 }
435 return retval;
436}
__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)
const int i
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
__global__ void const T *__restrict__ x
__global__ void const T *__restrict__ const T *__restrict__ dr
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
__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
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
static int neko_eb_env()
static int neko_tune_rounds()
#define NEKO_TUNE_LOG(LX, T1, T2)
#define NEKO_TUNE_INIT
#define NEKO_TUNE_BEST(T, BEST, N)
static int neko_tune_iters()
static int neko_chunks_env()
static int neko_eb_sweep()
#define NEKO_TUNE_WARMUP
#define NEKO_DMMA_CANDIDATES
static bool cuda_have_dmma()
static int neko_dmma_env()
#define NEKO_DMMA_PACK(LX)
#define NEKO_TUNE_LOG_DMMA(LX, T3)
#define NEKO_DMMA_NW(C)
static bool dmma_tma_cdtp_aligned(const void *dtx, const void *x, const void *dr, const void *ds, const void *dt)
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
static int neko_dmma_tma_env()
static bool cuda_have_tma()
void log_error(char *msg)
void log_message(char *msg)
void log_end_section()
void log_section(char *msg)
#define CASE_DMMA(LX, C)
#define CASE_DMMA_SEL(LX, SEL)
#define CASE_KSTEP_SEL(LX, SEL)
#define CASE(LX)
int tune_cdtp(void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
Definition opr_cdtp.cu:219
#define CASE_1D_SEL(LX, SEL)
#define CASE_KSTEP(LX, C)
#define CASE_LARGE(LX)
#define CASE_DMMA_TMA_SEL(LX, SEL)
void cuda_cdtp(void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx)
Definition opr_cdtp.cu:59
#define CASE_DMMA_TMA(LX, C)
#define CASE_1D(LX, C)